[PATCH 0/21] MR221: Draft: vkd3d: Implement tiled resources

[Conor McCarthy (＠cmccarthy)]

Depends on branch tiled_resource1, MR 216. The first new commit is "Add tests for GetResourceTiling()."

[PATCH 01/21] vkd3d: Check sparse image format is supported.

[Conor McCarthy]

From: Conor McCarthy

Based on code by Jan Sikorski.
---
 libs/vkd3d/resource.c | 16 ++++++++++++++++
 1 file changed, 16 insertions(+)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index ea7b6859c..e8b249bf5 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -776,9 +776,11 @@ static HRESULT vkd3d_create_image(struct d3d12_device *device,
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
     const struct vkd3d_format_compatibility_list *compat_list;
     const bool sparse_resource = !heap_properties;
+    VkSparseImageFormatProperties properties[2];
     VkImageFormatListCreateInfoKHR format_list;
     const struct vkd3d_format *format;
     VkImageCreateInfo image_info;
+    uint32_t count;
     VkResult vr;

if (resource)
@@ -914,6 +916,20 @@ static HRESULT vkd3d_create_image(struct d3d12_device *device,
     if (resource && image_info.tiling == VK_IMAGE_TILING_LINEAR)
         resource->flags |= VKD3D_RESOURCE_LINEAR_TILING;

+    if (sparse_resource)
+    {
+        count = ARRAY_SIZE(properties);
+        VK_CALL(vkGetPhysicalDeviceSparseImageFormatProperties(device->vk_physical_device, image_info.format,
+                image_info.imageType, image_info.samples, image_info.usage, image_info.tiling, &count, properties));
+
+        if (!count)
+        {
+            FIXME("Sparse images are not supported with format %u, type %u, samples %u, usage %#x.\n",
+                    image_info.format, image_info.imageType, image_info.samples, image_info.usage);
+            return E_INVALIDARG;
+        }
+    }
+
     if ((vr = VK_CALL(vkCreateImage(device->vk_device, &image_info, NULL, vk_image))) < 0)
         WARN("Failed to create Vulkan image, vr %d.\n", vr);

[PATCH 02/21] vkd3d: Always pass null heap properties to vkd3d_create_image() for sparse images.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/resource.c | 5 ++++-
 1 file changed, 4 insertions(+), 1 deletion(-)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index e8b249bf5..4b1b6044a 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -944,6 +944,7 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
     D3D12_RESOURCE_DESC validated_desc;
     VkMemoryRequirements requirements;
     VkImage vk_image;
+    bool tiled;
     HRESULT hr;

assert(desc->Dimension != D3D12_RESOURCE_DIMENSION_BUFFER);
@@ -956,8 +957,10 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
         desc = &validated_desc;
     }

+    tiled = desc->Layout == D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+
     /* XXX: We have to create an image to get its memory requirements. */
-    if (SUCCEEDED(hr = vkd3d_create_image(device, &heap_properties, 0, desc, NULL, &vk_image)))
+    if (SUCCEEDED(hr = vkd3d_create_image(device, tiled ? NULL : &heap_properties, 0, desc, NULL, &vk_image)))
     {
         VK_CALL(vkGetImageMemoryRequirements(device->vk_device, vk_image, &requirements));
         VK_CALL(vkDestroyImage(device->vk_device, vk_image, NULL));

[PATCH 03/21] vkd3d: Validate tiled resources tier for 3D textures.

[Conor McCarthy]

From: Conor McCarthy

Check directly for Vulkan support because the D3D12 tiled resources
tier may in future be modified by a config option.
---
 libs/vkd3d/device.c        | 1 +
 libs/vkd3d/resource.c      | 9 +++++++++
 libs/vkd3d/vkd3d_private.h | 4 +++-
 3 files changed, 13 insertions(+), 1 deletion(-)

diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 4263dcf41..007a6f651 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -1464,6 +1464,7 @@ static HRESULT vkd3d_init_device_caps(struct d3d12_device *device,

vulkan_info->device_limits = physical_device_info->properties2.properties.limits;
     vulkan_info->sparse_properties = physical_device_info->properties2.properties.sparseProperties;
+    vulkan_info->sparse_residency_3d = features->sparseResidencyImage3D;
     vulkan_info->rasterization_stream = physical_device_info->xfb_properties.transformFeedbackRasterizationStreamSelect;
     vulkan_info->transform_feedback_queries = physical_device_info->xfb_properties.transformFeedbackQueries;
     vulkan_info->uav_read_without_format = features->shaderStorageImageReadWithoutFormat;
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index 4b1b6044a..ebd020c12 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1680,6 +1680,15 @@ HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc, struct d3d
                 return E_INVALIDARG;
             }

+            if (desc->Layout == D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE)
+            {
+                if (desc->Dimension == D3D12_RESOURCE_DIMENSION_TEXTURE3D && !device->vk_info.sparse_residency_3d)
+                {
+                    WARN("The device does not support tiled 3D images.\n");
+                    return E_INVALIDARG;
+                }
+            }
+
             if (!d3d12_resource_validate_texture_format(desc, format)
                     || !d3d12_resource_validate_texture_alignment(desc, format))
                 return E_INVALIDARG;
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index b01507544..dde82414f 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -147,9 +147,11 @@ struct vkd3d_vulkan_info
     unsigned int max_vertex_attrib_divisor;

VkPhysicalDeviceLimits device_limits;
-    VkPhysicalDeviceSparseProperties sparse_properties;
     struct vkd3d_device_descriptor_limits descriptor_limits;

+    VkPhysicalDeviceSparseProperties sparse_properties;
+    bool sparse_residency_3d;
+
     VkPhysicalDeviceTexelBufferAlignmentPropertiesEXT texel_buffer_alignment_properties;

unsigned int shader_extension_count;

[PATCH 04/21] vkd3d: Validate plane count for tiled textures.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/resource.c | 6 ++++++
 1 file changed, 6 insertions(+)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index ebd020c12..e1c295922 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1687,6 +1687,12 @@ HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc, struct d3d
                     WARN("The device does not support tiled 3D images.\n");
                     return E_INVALIDARG;
                 }
+                if (format->plane_count > 1)
+                {
+                    WARN("Invalid format %#x. D3D12 does not support multiplanar formats for tiled resources.\n",
+                            format->dxgi_format);
+                    return E_INVALIDARG;
+                }
             }

if (!d3d12_resource_validate_texture_format(desc, format)

[PATCH 05/21] vkd3d: Add missing const attributes to ID3D12CommandQueue::UpdateTileMappings() parameters.

[Conor McCarthy]

From: Conor McCarthy

---
 include/vkd3d_d3d12.idl | 4 ++--
 libs/vkd3d/command.c    | 2 +-
 2 files changed, 3 insertions(+), 3 deletions(-)

diff --git a/include/vkd3d_d3d12.idl b/include/vkd3d_d3d12.idl
index 652ffc364..06287f5a8 100644
--- a/include/vkd3d_d3d12.idl
+++ b/include/vkd3d_d3d12.idl
@@ -2266,8 +2266,8 @@ interface ID3D12CommandQueue : ID3D12Pageable
             ID3D12Heap *heap,
             UINT range_count,
             const D3D12_TILE_RANGE_FLAGS *range_flags,
-            UINT *heap_range_offsets,
-            UINT *range_tile_counts,
+            const UINT *heap_range_offsets,
+            const UINT *range_tile_counts,
             D3D12_TILE_MAPPING_FLAGS flags);

void CopyTileMappings(ID3D12Resource *dst_resource,
diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 32439eec7..c5bd687bd 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -6166,7 +6166,7 @@ static void STDMETHODCALLTYPE d3d12_command_queue_UpdateTileMappings(ID3D12Comma
         ID3D12Resource *resource, UINT region_count,
         const D3D12_TILED_RESOURCE_COORDINATE *region_start_coordinates, const D3D12_TILE_REGION_SIZE *region_sizes,
         ID3D12Heap *heap, UINT range_count, const D3D12_TILE_RANGE_FLAGS *range_flags,
-        UINT *heap_range_offsets, UINT *range_tile_counts, D3D12_TILE_MAPPING_FLAGS flags)
+        const UINT *heap_range_offsets, const UINT *range_tile_counts, D3D12_TILE_MAPPING_FLAGS flags)
 {
     FIXME("iface %p, resource %p, region_count %u, region_start_coordinates %p, "
             "region_sizes %p, heap %p, range_count %u, range_flags %p, heap_range_offsets %p, "

[PATCH 06/21] vkd3d: Append UpdateTileMappings() commands to the command queue op array.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 114 ++++++++++++++++++++++++++++++++++++-
 libs/vkd3d/resource.c      |   9 +++
 libs/vkd3d/vkd3d_private.h |  25 ++++++++
 3 files changed, 146 insertions(+), 2 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index c5bd687bd..f7cec81ec 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -26,6 +26,7 @@ static HRESULT d3d12_fence_signal(struct d3d12_fence *fence, uint64_t value, VkF
 static void d3d12_fence_signal_timeline_semaphore(struct d3d12_fence *fence, uint64_t timeline_value);
 static HRESULT d3d12_command_queue_signal(struct d3d12_command_queue *command_queue,
         struct d3d12_fence *fence, uint64_t value);
+static void d3d12_command_queue_submit_locked(struct d3d12_command_queue *queue);
 static HRESULT d3d12_command_queue_flush_ops(struct d3d12_command_queue *queue, bool *flushed_any);
 static HRESULT d3d12_command_queue_flush_ops_locked(struct d3d12_command_queue *queue, bool *flushed_any);

@@ -6162,17 +6163,115 @@ static struct vkd3d_cs_op_data *d3d12_command_queue_op_array_require_space(struc
     return &array->ops[array->count++];
 }

+static bool clone_tile_mapping_parameter(void **dst, const void *src, size_t elem_size, unsigned int count)
+{
+    void *buffer;
+
+    *dst = NULL;
+    if (src)
+    {
+        if (!(buffer = vkd3d_malloc(count * elem_size)))
+            return false;
+        memcpy(buffer, src, count * elem_size);
+        *dst = buffer;
+    }
+    return true;
+}
+
 static void STDMETHODCALLTYPE d3d12_command_queue_UpdateTileMappings(ID3D12CommandQueue *iface,
         ID3D12Resource *resource, UINT region_count,
         const D3D12_TILED_RESOURCE_COORDINATE *region_start_coordinates, const D3D12_TILE_REGION_SIZE *region_sizes,
         ID3D12Heap *heap, UINT range_count, const D3D12_TILE_RANGE_FLAGS *range_flags,
         const UINT *heap_range_offsets, const UINT *range_tile_counts, D3D12_TILE_MAPPING_FLAGS flags)
 {
-    FIXME("iface %p, resource %p, region_count %u, region_start_coordinates %p, "
+    struct d3d12_resource *resource_impl = unsafe_impl_from_ID3D12Resource(resource);
+    struct d3d12_command_queue *command_queue = impl_from_ID3D12CommandQueue(iface);
+    struct d3d12_heap *heap_impl = unsafe_impl_from_ID3D12Heap(heap);
+    struct vkd3d_cs_op_data *op;
+
+    TRACE("iface %p, resource %p, region_count %u, region_start_coordinates %p, "
             "region_sizes %p, heap %p, range_count %u, range_flags %p, heap_range_offsets %p, "
-            "range_tile_counts %p, flags %#x stub!\n",
+            "range_tile_counts %p, flags %#x.\n",
             iface, resource, region_count, region_start_coordinates, region_sizes, heap, range_count,
             range_flags, heap_range_offsets, range_tile_counts, flags);
+
+    if (!region_count || !range_count)
+        return;
+
+    if (!command_queue->supports_sparse_binding)
+    {
+        FIXME("Command queue %p does not support sparse binding.\n", command_queue);
+        return;
+    }
+
+    if (!resource_impl->tiles.subresource_count)
+    {
+        WARN("Resource %p is not a tiled resource.\n", resource_impl);
+        return;
+    }
+
+    if (region_count > 1 && !region_start_coordinates)
+    {
+        WARN("Region start coordinates must not be NULL when region count is > 1.\n");
+        return;
+    }
+
+    if (range_count > 1 && !range_tile_counts)
+    {
+        WARN("Range tile counts must not be NULL when range count is > 1.\n");
+        return;
+    }
+
+    vkd3d_mutex_lock(&command_queue->op_mutex);
+
+    if (!(op = d3d12_command_queue_op_array_require_space(&command_queue->op_queue)))
+    {
+        ERR("Failed to add op.\n");
+        goto done;
+    }
+
+    op->opcode = VKD3D_CS_OP_UPDATE_MAPPINGS;
+    memset(&op->u.update_mappings, 0, sizeof(op->u.update_mappings));
+    op->u.update_mappings.resource = resource_impl;
+    op->u.update_mappings.heap = heap_impl;
+    if (!clone_tile_mapping_parameter((void **)&op->u.update_mappings.region_start_coordinates,
+            region_start_coordinates, sizeof(*region_start_coordinates), region_count))
+    {
+        ERR("Failed to allocate region start coordinates.\n");
+        goto done;
+    }
+    if (!clone_tile_mapping_parameter((void **)&op->u.update_mappings.region_sizes,
+            region_sizes, sizeof(*region_sizes), region_count))
+    {
+        ERR("Failed to allocate region sizes.\n");
+        goto done;
+    }
+    if (!clone_tile_mapping_parameter((void **)&op->u.update_mappings.range_flags,
+            range_flags, sizeof(*range_flags), range_count))
+    {
+        ERR("Failed to allocate range flags.\n");
+        goto done;
+    }
+    if (!clone_tile_mapping_parameter((void **)&op->u.update_mappings.heap_range_offsets,
+            heap_range_offsets, sizeof(*heap_range_offsets), range_count))
+    {
+        ERR("Failed to allocate heap range offsets.\n");
+        goto done;
+    }
+    if (!clone_tile_mapping_parameter((void **)&op->u.update_mappings.range_tile_counts,
+            range_tile_counts, sizeof(*range_tile_counts), range_count))
+    {
+        ERR("Failed to allocate range tile counts.\n");
+        goto done;
+    }
+    op->u.update_mappings.region_count = region_count;
+    op->u.update_mappings.range_count = range_count;
+    op->u.update_mappings.flags = flags;
+
+    d3d12_command_queue_submit_locked(command_queue);
+
+done:
+    vkd3d_mutex_unlock(&command_queue->op_mutex);
 }

static void STDMETHODCALLTYPE d3d12_command_queue_CopyTileMappings(ID3D12CommandQueue *iface,
@@ -6934,6 +7033,15 @@ static HRESULT d3d12_command_queue_flush_ops_locked(struct d3d12_command_queue *
                     d3d12_command_queue_execute(queue, op->u.execute.buffers, op->u.execute.buffer_count);
                     break;

+                case VKD3D_CS_OP_UPDATE_MAPPINGS:
+                    FIXME("Tiled resource binding is not supported yet.\n");
+                    vkd3d_free(op->u.update_mappings.region_start_coordinates);
+                    vkd3d_free(op->u.update_mappings.region_sizes);
+                    vkd3d_free(op->u.update_mappings.range_flags);
+                    vkd3d_free(op->u.update_mappings.heap_range_offsets);
+                    vkd3d_free(op->u.update_mappings.range_tile_counts);
+                    break;
+
                 default:
                     vkd3d_unreachable();
             }
@@ -7000,6 +7108,8 @@ static HRESULT d3d12_command_queue_init(struct d3d12_command_queue *queue,
     if (FAILED(hr = vkd3d_fence_worker_start(&queue->fence_worker, queue->vkd3d_queue, device)))
         goto fail_destroy_op_mutex;

+    queue->supports_sparse_binding = !!(queue->vkd3d_queue->vk_queue_flags & VK_QUEUE_SPARSE_BINDING_BIT);
+
     d3d12_device_add_ref(queue->device = device);

return S_OK;
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index e1c295922..709de8ab2 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1058,6 +1058,11 @@ static void d3d12_resource_get_level_box(const struct d3d12_resource *resource,
     box->back = d3d12_resource_desc_get_depth(&resource->desc, level);
 }

+static void d3d12_resource_init_tiles(struct d3d12_resource *resource)
+{
+    resource->tiles.subresource_count = d3d12_resource_desc_get_sub_resource_count(&resource->desc);
+}
+
 /* ID3D12Resource */
 static inline struct d3d12_resource *impl_from_ID3D12Resource(ID3D12Resource *iface)
 {
@@ -1821,6 +1826,8 @@ static HRESULT d3d12_resource_init(struct d3d12_resource *resource, struct d3d12
     resource->heap = NULL;
     resource->heap_offset = 0;

+    memset(&resource->tiles, 0, sizeof(resource->tiles));
+
     if (FAILED(hr = vkd3d_private_store_init(&resource->private_store)))
     {
         d3d12_resource_destroy(resource, device);
@@ -2006,6 +2013,8 @@ HRESULT d3d12_reserved_resource_create(struct d3d12_device *device,
             desc, initial_state, optimized_clear_value, &object)))
         return hr;

+    d3d12_resource_init_tiles(object);
+
     TRACE("Created reserved resource %p.\n", object);

*resource = object;
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index dde82414f..b4a544828 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -672,6 +672,11 @@ struct d3d12_heap *unsafe_impl_from_ID3D12Heap(ID3D12Heap *iface);
 #define VKD3D_RESOURCE_DEDICATED_HEAP 0x00000008
 #define VKD3D_RESOURCE_LINEAR_TILING  0x00000010

+struct d3d12_resource_tile_info
+{
+    unsigned int subresource_count;
+};
+
 /* ID3D12Resource */
 struct d3d12_resource
 {
@@ -700,6 +705,8 @@ struct d3d12_resource

struct d3d12_device *device;

+    struct d3d12_resource_tile_info tiles;
+
     struct vkd3d_private_store private_store;
 };

@@ -1456,6 +1463,7 @@ enum vkd3d_cs_op
     VKD3D_CS_OP_WAIT,
     VKD3D_CS_OP_SIGNAL,
     VKD3D_CS_OP_EXECUTE,
+    VKD3D_CS_OP_UPDATE_MAPPINGS,
 };

struct vkd3d_cs_wait
@@ -1476,6 +1484,20 @@ struct vkd3d_cs_execute
     unsigned int buffer_count;
 };

+struct vkd3d_cs_update_mappings
+{
+    struct d3d12_resource *resource;
+    struct d3d12_heap *heap;
+    D3D12_TILED_RESOURCE_COORDINATE *region_start_coordinates;
+    D3D12_TILE_REGION_SIZE *region_sizes;
+    D3D12_TILE_RANGE_FLAGS *range_flags;
+    UINT *heap_range_offsets;
+    UINT *range_tile_counts;
+    UINT region_count;
+    UINT range_count;
+    D3D12_TILE_MAPPING_FLAGS flags;
+};
+
 struct vkd3d_cs_op_data
 {
     enum vkd3d_cs_op opcode;
@@ -1484,6 +1506,7 @@ struct vkd3d_cs_op_data
         struct vkd3d_cs_wait wait;
         struct vkd3d_cs_signal signal;
         struct vkd3d_cs_execute execute;
+        struct vkd3d_cs_update_mappings update_mappings;
     } u;
 };

@@ -1521,6 +1544,8 @@ struct d3d12_command_queue
      * set, aux_op_queue.count must be zero. */
     struct d3d12_command_queue_op_array aux_op_queue;

+    bool supports_sparse_binding;
+
     struct vkd3d_private_store private_store;
 };

[PATCH 07/21] vkd3d: Append CopyTileMappings() commands to the command queue op array.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 32 ++++++++++++++++++++++++++++++--
 libs/vkd3d/resource.c      |  5 -----
 libs/vkd3d/vkd3d_private.h | 17 +++++++++++++++++
 3 files changed, 47 insertions(+), 7 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index f7cec81ec..39c0f4cd3 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -6282,10 +6282,34 @@ static void STDMETHODCALLTYPE d3d12_command_queue_CopyTileMappings(ID3D12Command
         const D3D12_TILE_REGION_SIZE *region_size,
         D3D12_TILE_MAPPING_FLAGS flags)
 {
-    FIXME("iface %p, dst_resource %p, dst_region_start_coordinate %p, "
-            "src_resource %p, src_region_start_coordinate %p, region_size %p, flags %#x stub!\n",
+    struct d3d12_resource *dst_resource_impl = impl_from_ID3D12Resource(dst_resource);
+    struct d3d12_resource *src_resource_impl = impl_from_ID3D12Resource(src_resource);
+    struct d3d12_command_queue *command_queue = impl_from_ID3D12CommandQueue(iface);
+    struct vkd3d_cs_op_data *op;
+
+    TRACE("iface %p, dst_resource %p, dst_region_start_coordinate %p, "
+            "src_resource %p, src_region_start_coordinate %p, region_size %p, flags %#x.\n",
             iface, dst_resource, dst_region_start_coordinate, src_resource,
             src_region_start_coordinate, region_size, flags);
+
+    vkd3d_mutex_lock(&command_queue->op_mutex);
+
+    if (!(op = d3d12_command_queue_op_array_require_space(&command_queue->op_queue)))
+    {
+        ERR("Failed to add op.\n");
+        return;
+    }
+    op->opcode = VKD3D_CS_OP_COPY_MAPPINGS;
+    op->u.copy_mappings.dst_resource = dst_resource_impl;
+    op->u.copy_mappings.src_resource = src_resource_impl;
+    op->u.copy_mappings.dst_region_start_coordinate = *dst_region_start_coordinate;
+    op->u.copy_mappings.src_region_start_coordinate = *src_region_start_coordinate;
+    op->u.copy_mappings.region_size = *region_size;
+    op->u.copy_mappings.flags = flags;
+
+    d3d12_command_queue_submit_locked(command_queue);
+
+    vkd3d_mutex_unlock(&command_queue->op_mutex);
 }

static void d3d12_command_queue_execute(struct d3d12_command_queue *command_queue,
@@ -7042,6 +7066,10 @@ static HRESULT d3d12_command_queue_flush_ops_locked(struct d3d12_command_queue *
                     vkd3d_free(op->u.update_mappings.range_tile_counts);
                     break;

+                case VKD3D_CS_OP_COPY_MAPPINGS:
+                    FIXME("Tiled resource mapping copying is not supported yet.\n");
+                    break;
+
                 default:
                     vkd3d_unreachable();
             }
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index 709de8ab2..3ca2ff98f 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1064,11 +1064,6 @@ static void d3d12_resource_init_tiles(struct d3d12_resource *resource)
 }

/* ID3D12Resource */
-static inline struct d3d12_resource *impl_from_ID3D12Resource(ID3D12Resource *iface)
-{
-    return CONTAINING_RECORD(iface, struct d3d12_resource, ID3D12Resource_iface);
-}
-
 static HRESULT STDMETHODCALLTYPE d3d12_resource_QueryInterface(ID3D12Resource *iface,
         REFIID riid, void **object)
 {
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index b4a544828..48d66ad4a 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -710,6 +710,11 @@ struct d3d12_resource
     struct vkd3d_private_store private_store;
 };

+static inline struct d3d12_resource *impl_from_ID3D12Resource(ID3D12Resource *iface)
+{
+    return CONTAINING_RECORD(iface, struct d3d12_resource, ID3D12Resource_iface);
+}
+
 static inline bool d3d12_resource_is_buffer(const struct d3d12_resource *resource)
 {
     return resource->desc.Dimension == D3D12_RESOURCE_DIMENSION_BUFFER;
@@ -1464,6 +1469,7 @@ enum vkd3d_cs_op
     VKD3D_CS_OP_SIGNAL,
     VKD3D_CS_OP_EXECUTE,
     VKD3D_CS_OP_UPDATE_MAPPINGS,
+    VKD3D_CS_OP_COPY_MAPPINGS,
 };

struct vkd3d_cs_wait
@@ -1498,6 +1504,16 @@ struct vkd3d_cs_update_mappings
     D3D12_TILE_MAPPING_FLAGS flags;
 };

+struct vkd3d_cs_copy_mappings
+{
+    struct d3d12_resource *dst_resource;
+    struct d3d12_resource *src_resource;
+    D3D12_TILED_RESOURCE_COORDINATE dst_region_start_coordinate;
+    D3D12_TILED_RESOURCE_COORDINATE src_region_start_coordinate;
+    D3D12_TILE_REGION_SIZE region_size;
+    D3D12_TILE_MAPPING_FLAGS flags;
+};
+
 struct vkd3d_cs_op_data
 {
     enum vkd3d_cs_op opcode;
@@ -1507,6 +1523,7 @@ struct vkd3d_cs_op_data
         struct vkd3d_cs_signal signal;
         struct vkd3d_cs_execute execute;
         struct vkd3d_cs_update_mappings update_mappings;
+        struct vkd3d_cs_copy_mappings copy_mappings;
     } u;
 };

[PATCH 08/21] tests/d3d12: Add tests for GetResourceTiling().

[Conor McCarthy]

From: Conor McCarthy

Based on vkd3d-proton patches by Philip Rebohle and
Hans-Kristian Arntzen.
---
 tests/d3d12.c | 275 ++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 275 insertions(+)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 2f1c905f2..9f00d74fb 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -36438,6 +36438,280 @@ static void test_vs_ps_relative_addressing(void)
     destroy_test_context(&context);
 }

+static uint32_t compute_tile_count(uint32_t resource_size, uint32_t mip, uint32_t tile_size)
+{
+    uint32_t mip_size = max(resource_size >> mip, 1u);
+    return (mip_size + tile_size - 1) / tile_size;
+}
+
+static void test_get_resource_tiling(void)
+{
+    D3D12_SUBRESOURCE_TILING tilings_alt[17];
+    D3D12_PACKED_MIP_INFO packed_mip_info;
+    D3D12_SUBRESOURCE_TILING tilings[17];
+    UINT num_resource_tiles, num_tilings;
+    D3D12_RESOURCE_DESC resource_desc;
+    struct test_context_desc desc;
+    struct test_context context;
+    D3D12_TILE_SHAPE tile_shape;
+    ID3D12Resource *resource;
+    unsigned int i, j;
+    bool no_tier_3;
+    HRESULT hr;
+
+    static const struct
+    {
+        D3D12_RESOURCE_DIMENSION dim;
+        DXGI_FORMAT format;
+        UINT width;
+        UINT height;
+        UINT depth_or_array_layers;
+        UINT mip_levels;
+        UINT expected_tile_count;
+        UINT expected_tiling_count;
+        UINT expected_standard_mips;
+        UINT tile_shape_w;
+        UINT tile_shape_h;
+        UINT tile_shape_d;
+        D3D12_TILED_RESOURCES_TIER min_tier;
+        bool todo_radv;
+    }
+    tests[] =
+    {
+        /* Test buffers */
+        { D3D12_RESOURCE_DIMENSION_BUFFER,    DXGI_FORMAT_UNKNOWN,            1024,    1,  1,  1,  1,  1,  0, 65536,   1,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_BUFFER,    DXGI_FORMAT_UNKNOWN,        16*65536,    1,  1,  1, 16,  1,  0, 65536,   1,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Test small resource behavior */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,              1,    1,  1,  1,  1,  1,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,              2,    2,  1,  2,  1,  2,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,              4,    4,  1,  3,  1,  3,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,              8,    8,  1,  4,  1,  4,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,             16,   16,  1,  5,  1,  5,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,             32,   32,  1,  6,  1,  6,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,             64,   64,  1,  7,  1,  7,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,            128,  128,  1,  8,  1,  8,  0,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,            256,  256,  1,  9,  2,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Test various image formats */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8_UNORM,	           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8_UNORM,          512,  512,  1,  1,  8,  1,  1,   256, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      512,  512,  1,  1, 16,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R16G16B16A16_UNORM,  512,  512,  1,  1, 32,  1,  1,   128,  64,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R32G32B32A32_FLOAT,  512,  512,  1,  1, 64,  1,  1,    64,  64,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_D16_UNORM,           512,  512,  1,  1,  8,  1,  1,   256, 128,   1, D3D12_TILED_RESOURCES_TIER_1, true },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_D32_FLOAT,           512,  512,  1,  1, 16,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1, true },
+        /* Test rectangular textures */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,     1024,  256,  1,  1, 16,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      256, 1024,  1,  1, 16,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      192,  128,  1,  1,  2,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_2 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      128,  192,  1,  1,  2,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_2 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      320,  192,  1,  1,  6,  1,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_2 },
+        /* Test array layers and packed mip levels */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      128,  128, 16,  1, 16, 16,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      128,  128,  1,  8,  1,  8,  1,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      512,  512,  1, 10, 21, 10,  3,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,      512,  512,  4,  3, 84, 12,  3,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_R8G8B8A8_UNORM,       64,   64,  1,  1,  0,  1,  0,   128, 128,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Test 3D textures */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE3D, DXGI_FORMAT_R8_UNORM,             64,   64, 64,  1,  4,  1,  1,    64,  32,  32, D3D12_TILED_RESOURCES_TIER_3 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE3D, DXGI_FORMAT_R8G8_UNORM,           64,   64, 64,  1,  8,  1,  1,    32,  32,  32, D3D12_TILED_RESOURCES_TIER_3 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE3D, DXGI_FORMAT_R8G8B8A8_UNORM,       64,   64, 64,  1, 16,  1,  1,    32,  32,  16, D3D12_TILED_RESOURCES_TIER_3 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE3D, DXGI_FORMAT_R32G32B32A32_FLOAT,   64,   64, 64,  3, 73,  3,  3,    16,  16,  16, D3D12_TILED_RESOURCES_TIER_3 },
+        /* Basic BC configurations. */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,           512,  512,  1,  1,  2,  1,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC2_UNORM,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC3_UNORM,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC4_UNORM,           512,  512,  1,  1,  2,  1,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC5_UNORM,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC6H_UF16,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC6H_SF16,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           512,  512,  1,  1,  4,  1,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Basic mipmapping with obvious tiling layouts. */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,           512,  256,  1,  10,  2, 10,  1,  512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,          1024,  512,  1,  10,  6, 10,  2,  512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,          2048, 1024,  1,  10, 22, 10,  3,  512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           256,  256,  1,   9,  2,  9,  1,  256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           512,  512,  1,   9,  6,  9,  2,  256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,          1024, 1024,  1,   9, 22,  9,  3,  256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Wide shapes. On AMD, we keep observing standard mips even when the smallest dimension dips below the tile size.
+         * This is not the case on NV however. */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,          1024,  256,  1, 10,  3, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,          2048,  256,  1, 10,  6, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,          4096,  256,  1, 10, 11, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           512,  256,  1,  9,  3,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,          1024,  256,  1,  9,  6,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,          2048,  256,  1,  9, 11,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        /* Tall shapes. Similar to wide tests. */
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,           512,  512,  1, 10,  3, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,           512, 1024,  1, 10,  6, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC1_UNORM,           512, 2048,  1, 10, 11, 10,  1,   512, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           256,  512,  1,  9,  3,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           256, 1024,  1,  9,  6,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+        { D3D12_RESOURCE_DIMENSION_TEXTURE2D, DXGI_FORMAT_BC7_UNORM,           256, 2048,  1,  9, 11,  9,  1,   256, 256,   1, D3D12_TILED_RESOURCES_TIER_1 },
+    };
+
+    memset(&desc, 0, sizeof(desc));
+    desc.rt_width = 640;
+    desc.rt_height = 480;
+    desc.rt_format = DXGI_FORMAT_R8G8B8A8_UNORM;
+    if (!init_test_context(&context, &desc))
+        return;
+
+    /* Test behaviour with various parameter combinations */
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
+    resource_desc.Alignment = 0;
+    resource_desc.Width = 512;
+    resource_desc.Height = 512;
+    resource_desc.DepthOrArraySize = 1;
+    resource_desc.MipLevels = 10;
+    resource_desc.Format = DXGI_FORMAT_R8G8B8A8_UNORM;
+    resource_desc.SampleDesc.Count = 1;
+    resource_desc.SampleDesc.Quality = 0;
+    resource_desc.Layout = D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+    resource_desc.Flags = D3D12_RESOURCE_FLAG_NONE;
+
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_GENERIC_READ, NULL, &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == S_OK, "Failed to create reserved resource, hr %#x.\n", hr);
+
+    /* This is nonsense, but it doesn't crash or generate errors. */
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, NULL, 0, NULL);
+
+    /* If num_tilings is NULL, tilings_alt is ignored. */
+    memset(tilings, 0, sizeof(tilings));
+    memset(tilings_alt, 0, sizeof(tilings_alt));
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, NULL, 0, tilings_alt);
+    ok(memcmp(tilings, tilings_alt, sizeof(tilings_alt)) == 0, "Mismatch.\n");
+
+    num_tilings = 0;
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 0, NULL);
+    ok(num_tilings == 0, "Unexpected tiling count %u.\n", num_tilings);
+
+    num_tilings = ARRAY_SIZE(tilings);
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 10, tilings);
+    todo ok(num_tilings == 0, "Unexpected tiling count %u.\n", num_tilings);
+
+    num_tilings = ARRAY_SIZE(tilings);
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 2, tilings);
+    todo ok(num_tilings == 8, "Unexpected tiling count %u.\n", num_tilings);
+    todo ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);
+
+    num_tilings = 1;
+    memset(&tilings, 0xaa, sizeof(tilings));
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 2, tilings);
+    ok(num_tilings == 1, "Unexpected tiling count %u.\n", num_tilings);
+    todo ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);
+    ok(tilings[1].StartTileIndexInOverallResource == 0xaaaaaaaa, "Tiling array got modified.\n");
+
+    ID3D12Resource_Release(resource);
+
+    /* Tiled tier is not included in feature support yet so as not to break future bisections. */
+    no_tier_3 = is_amd_windows_device(context.device) || is_radv_device(context.device);
+
+    /* Test actual tiling properties */
+    for (i = 0; i < ARRAY_SIZE(tests); i++, vkd3d_test_pop_context())
+    {
+        unsigned int tile_index = 0;
+        vkd3d_test_push_context("test %u", i);
+
+        if (no_tier_3 && tests[i].min_tier > D3D12_TILED_RESOURCES_TIER_2)
+        {
+            skip("Tiled resources tier %u not supported.\n", tests[i].min_tier);
+            continue;
+        }
+
+        memset(&packed_mip_info, 0xaa, sizeof(packed_mip_info));
+        memset(&tile_shape, 0xaa, sizeof(tile_shape));
+        memset(&tilings, 0xaa, sizeof(tilings));
+
+        num_resource_tiles = 0xdeadbeef;
+        num_tilings = ARRAY_SIZE(tilings);
+
+        resource_desc.Dimension = tests[i].dim;
+        resource_desc.Alignment = 0;
+        resource_desc.Width = tests[i].width;
+        resource_desc.Height = tests[i].height;
+        resource_desc.DepthOrArraySize = tests[i].depth_or_array_layers;
+        resource_desc.MipLevels = tests[i].mip_levels;
+        resource_desc.Format = tests[i].format;
+        resource_desc.SampleDesc.Count = 1;
+        resource_desc.SampleDesc.Quality = 0;
+        resource_desc.Layout = D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+        resource_desc.Flags = D3D12_RESOURCE_FLAG_NONE;
+
+        if (tests[i].dim == D3D12_RESOURCE_DIMENSION_BUFFER)
+            resource_desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+
+        hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+                D3D12_RESOURCE_STATE_GENERIC_READ, NULL, &IID_ID3D12Resource, (void **)&resource);
+        todo_if(is_radv_device(context.device) && tests[i].todo_radv)
+        ok(hr == S_OK, "Failed to create reserved resource, hr %#x.\n", hr);
+
+        if (hr != S_OK)
+            continue;
+
+        ID3D12Device_GetResourceTiling(context.device, resource, &num_resource_tiles, &packed_mip_info, &tile_shape, &num_tilings, 0, tilings);
+
+        todo ok(num_resource_tiles != 0xdeadbeef && num_resource_tiles >= tests[i].expected_tile_count,
+                "Unexpected resource tile count %u.\n", num_resource_tiles);
+        todo ok(num_tilings == tests[i].expected_tiling_count, "Unexpected subresource tiling count %u.\n", num_tilings);
+
+        todo ok(packed_mip_info.NumStandardMips != 0xaa && packed_mip_info.NumStandardMips >= tests[i].expected_standard_mips,
+                "Unexpected standard mip count %u.\n", packed_mip_info.NumStandardMips);
+        todo ok(packed_mip_info.NumPackedMips == (tests[i].dim == D3D12_RESOURCE_DIMENSION_BUFFER
+                ? 0 : tests[i].mip_levels - packed_mip_info.NumStandardMips),
+                "Unexpected packed mip count %u.\n", packed_mip_info.NumPackedMips);
+        todo ok(packed_mip_info.NumPackedMips != 0xaa
+                && (packed_mip_info.NumTilesForPackedMips == 0) == (packed_mip_info.NumPackedMips == 0),
+                "Unexpected packed tile count %u.\n", packed_mip_info.NumTilesForPackedMips);
+
+        /* Docs say that tile shape should be cleared to zero if there is no standard mip, but drivers don't seem to care about this. */
+        todo ok(tile_shape.WidthInTexels == tests[i].tile_shape_w, "Unexpected tile width %u.\n", tile_shape.WidthInTexels);
+        todo ok(tile_shape.HeightInTexels == tests[i].tile_shape_h, "Unexpected tile height %u.\n", tile_shape.HeightInTexels);
+        todo ok(tile_shape.DepthInTexels == tests[i].tile_shape_d, "Unexpected tile depth %u.\n", tile_shape.DepthInTexels);
+
+        for (j = 0; j < tests[i].expected_tiling_count; j++)
+        {
+            uint32_t mip = j % tests[i].mip_levels;
+
+            if (mip < packed_mip_info.NumStandardMips || !packed_mip_info.NumPackedMips)
+            {
+                uint32_t expected_w = compute_tile_count(tests[i].width, mip, tests[i].tile_shape_w);
+                uint32_t expected_h = compute_tile_count(tests[i].height, mip, tests[i].tile_shape_h);
+                uint32_t expected_d = 1;
+
+                if (tests[i].dim == D3D12_RESOURCE_DIMENSION_TEXTURE3D)
+                    expected_d = compute_tile_count(tests[i].depth_or_array_layers, mip, tests[i].tile_shape_d);
+
+                todo ok(tilings[j].WidthInTiles == expected_w, "Unexpected width %u for subresource %u.\n", tilings[j].WidthInTiles, j);
+                todo ok(tilings[j].HeightInTiles == expected_h, "Unexpected width %u for subresource %u.\n", tilings[j].HeightInTiles, j);
+                todo ok(tilings[j].DepthInTiles == expected_d, "Unexpected width %u for subresource %u.\n", tilings[j].DepthInTiles, j);
+
+                todo ok(tilings[j].StartTileIndexInOverallResource == tile_index, "Unexpected start tile index %u for subresource %u.\n",
+                        tilings[j].StartTileIndexInOverallResource, j);
+
+                tile_index += tilings[j].WidthInTiles * tilings[j].HeightInTiles * tilings[j].DepthInTiles;
+            }
+            else
+            {
+                todo ok(!tilings[j].WidthInTiles && !tilings[j].HeightInTiles && !tilings[j].DepthInTiles,
+                        "Unexpected tile count (%u,%u,%u) for packed subresource %u.\n",
+                        tilings[j].WidthInTiles, tilings[j].HeightInTiles, tilings[j].DepthInTiles, j);
+                todo ok(tilings[j].StartTileIndexInOverallResource == 0xffffffff, "Unexpected start tile index %u for packed subresource %u.\n",
+                        tilings[j].StartTileIndexInOverallResource, j);
+            }
+        }
+
+        todo ok(num_resource_tiles == tile_index + packed_mip_info.NumTilesForPackedMips,
+                "Unexpected resource tile count %u.\n", num_resource_tiles);
+        todo ok(packed_mip_info.StartTileIndexInOverallResource == (packed_mip_info.NumPackedMips ? tile_index : 0),
+                "Unexpected mip tail start tile index %u.\n", packed_mip_info.StartTileIndexInOverallResource);
+
+        ID3D12Resource_Release(resource);
+    }
+
+    destroy_test_context(&context);
+}
+
 START_TEST(d3d12)
 {
     parse_args(argc, argv);
@@ -36617,4 +36891,5 @@ START_TEST(d3d12)
     run_test(test_clock_calibration);
     run_test(test_readback_map_stability);
     run_test(test_vs_ps_relative_addressing);
+    run_test(test_get_resource_tiling);
 }

[PATCH 09/21] vkd3d: Implement ID3D12Device::GetResourceTiling() for buffers.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/device.c        |  16 +++++-
 libs/vkd3d/resource.c      | 102 +++++++++++++++++++++++++++++++++++--
 libs/vkd3d/vkd3d_private.h |  24 +++++++++
 tests/d3d12.c              |  38 ++++++++------
 4 files changed, 159 insertions(+), 21 deletions(-)

diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 007a6f651..69dbb8705 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -3890,12 +3890,24 @@ static void STDMETHODCALLTYPE d3d12_device_GetResourceTiling(ID3D12Device *iface
         UINT *sub_resource_tiling_count, UINT first_sub_resource_tiling,
         D3D12_SUBRESOURCE_TILING *sub_resource_tilings)
 {
-    FIXME("iface %p, resource %p, total_tile_count %p, packed_mip_info %p, "
+    const struct d3d12_resource *resource_impl = impl_from_ID3D12Resource(resource);
+    struct d3d12_device *device = impl_from_ID3D12Device(iface);
+
+    TRACE("iface %p, resource %p, total_tile_count %p, packed_mip_info %p, "
             "standard_title_shape %p, sub_resource_tiling_count %p, "
-            "first_sub_resource_tiling %u, sub_resource_tilings %p stub!\n",
+            "first_sub_resource_tiling %u, sub_resource_tilings %p.\n",
             iface, resource, total_tile_count, packed_mip_info, standard_tile_shape,
             sub_resource_tiling_count, first_sub_resource_tiling,
             sub_resource_tilings);
+
+    if (d3d12_resource_is_texture(resource_impl))
+    {
+        FIXME("Not implemented for textures.\n");
+        return;
+    }
+
+    d3d12_resource_get_tiling(device, resource_impl, total_tile_count, packed_mip_info, standard_tile_shape,
+            sub_resource_tiling_count, first_sub_resource_tiling, sub_resource_tilings);
 }

static LUID * STDMETHODCALLTYPE d3d12_device_GetAdapterLuid(ID3D12Device *iface, LUID *luid)
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index 3ca2ff98f..c50d6241f 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -972,6 +972,11 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
     return hr;
 }

+static void d3d12_resource_tile_info_cleanup(struct d3d12_resource *resource)
+{
+    vkd3d_free(resource->tiles.subresources);
+}
+
 static void d3d12_resource_destroy(struct d3d12_resource *resource, struct d3d12_device *device)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
@@ -987,6 +992,8 @@ static void d3d12_resource_destroy(struct d3d12_resource *resource, struct d3d12
     else
         VK_CALL(vkDestroyImage(device->vk_device, resource->u.vk_image, NULL));

+    d3d12_resource_tile_info_cleanup(resource);
+
     if (resource->heap)
         d3d12_heap_resource_destroyed(resource->heap);
 }
@@ -1058,9 +1065,94 @@ static void d3d12_resource_get_level_box(const struct d3d12_resource *resource,
     box->back = d3d12_resource_desc_get_depth(&resource->desc, level);
 }

-static void d3d12_resource_init_tiles(struct d3d12_resource *resource)
+void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_resource *resource,
+        UINT *total_tile_count, D3D12_PACKED_MIP_INFO *packed_mip_info, D3D12_TILE_SHAPE *standard_tile_shape,
+        UINT *subresource_tiling_count, UINT first_subresource_tiling,
+        D3D12_SUBRESOURCE_TILING *subresource_tilings)
 {
-    resource->tiles.subresource_count = d3d12_resource_desc_get_sub_resource_count(&resource->desc);
+    unsigned int i, subresource, subresource_count, count;
+    const struct vkd3d_subresource_tile_info *tile_info;
+    const VkExtent3D *tile_extent;
+
+    tile_extent = &resource->tiles.tile_extent;
+
+    if (packed_mip_info)
+    {
+        packed_mip_info->NumStandardMips = resource->tiles.standard_mip_count;
+        packed_mip_info->NumPackedMips = 0;
+        packed_mip_info->NumTilesForPackedMips = 0;
+        packed_mip_info->StartTileIndexInOverallResource = 0;
+    }
+
+    if (standard_tile_shape)
+    {
+        /* D3D12 docs say tile shape is cleared to zero if there is no standard mip, but drivers don't seem to do this. */
+        standard_tile_shape->WidthInTexels = tile_extent->width;
+        standard_tile_shape->HeightInTexels = tile_extent->height;
+        standard_tile_shape->DepthInTexels = tile_extent->depth;
+    }
+
+    if (total_tile_count)
+        *total_tile_count = resource->tiles.total_count;
+
+    if (!subresource_tiling_count)
+        return;
+
+    subresource_count = resource->tiles.subresource_count;
+
+    count = subresource_count - min(first_subresource_tiling, subresource_count);
+    count = min(count, *subresource_tiling_count);
+
+    for (i = 0; i < count; ++i)
+    {
+        subresource = i + first_subresource_tiling;
+        tile_info = &resource->tiles.subresources[subresource];
+        subresource_tilings[i].StartTileIndexInOverallResource = tile_info->offset;
+        subresource_tilings[i].WidthInTiles = tile_info->extent.width;
+        subresource_tilings[i].HeightInTiles = tile_info->extent.height;
+        subresource_tilings[i].DepthInTiles = tile_info->extent.depth;
+    }
+    *subresource_tiling_count = i;
+}
+
+static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3d12_device *device)
+{
+    struct vkd3d_subresource_tile_info *tile_info;
+    unsigned int subresource_count;
+
+    subresource_count = d3d12_resource_desc_get_sub_resource_count(&resource->desc);
+
+    if (!d3d12_resource_is_buffer(resource))
+        return true;
+
+    if (!(resource->tiles.subresources = vkd3d_calloc(subresource_count, sizeof(*resource->tiles.subresources))))
+    {
+        ERR("Failed to allocate subresource info array.\n");
+        return false;
+    }
+
+    if (d3d12_resource_is_buffer(resource))
+    {
+        tile_info = &resource->tiles.subresources[0];
+        tile_info->offset = 0;
+        tile_info->extent.width = align(resource->desc.Width, D3D12_TILE_SIZE) / D3D12_TILE_SIZE;
+        tile_info->extent.height = 1;
+        tile_info->extent.depth = 1;
+        tile_info->count = tile_info->extent.width;
+
+        resource->tiles.tile_extent.width = D3D12_TILE_SIZE;
+        resource->tiles.tile_extent.height = 1;
+        resource->tiles.tile_extent.depth = 1;
+        resource->tiles.total_count = tile_info->extent.width;
+        resource->tiles.subresource_count = 1;
+        resource->tiles.standard_mip_count = 1;
+    }
+    else
+    {
+        vkd3d_unreachable();
+    }
+
+    return true;
 }

/* ID3D12Resource */
@@ -2008,7 +2100,11 @@ HRESULT d3d12_reserved_resource_create(struct d3d12_device *device,
             desc, initial_state, optimized_clear_value, &object)))
         return hr;

-    d3d12_resource_init_tiles(object);
+    if (!d3d12_resource_init_tiles(object, device))
+    {
+        d3d12_resource_Release(&object->ID3D12Resource_iface);
+        return E_OUTOFMEMORY;
+    }

TRACE("Created reserved resource %p.\n", object);

diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 48d66ad4a..b545945a5 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -672,11 +672,31 @@ struct d3d12_heap *unsafe_impl_from_ID3D12Heap(ID3D12Heap *iface);
 #define VKD3D_RESOURCE_DEDICATED_HEAP 0x00000008
 #define VKD3D_RESOURCE_LINEAR_TILING  0x00000010

+struct vkd3d_tiled_region_extent
+{
+    unsigned int width;
+    unsigned int height;
+    unsigned int depth;
+};
+
+struct vkd3d_subresource_tile_info
+{
+    unsigned int offset;
+    unsigned int count;
+    struct vkd3d_tiled_region_extent extent;
+};
+
 struct d3d12_resource_tile_info
 {
+    VkExtent3D tile_extent;
+    unsigned int total_count;
+    unsigned int standard_mip_count;
     unsigned int subresource_count;
+    struct vkd3d_subresource_tile_info *subresources;
 };

+#define D3D12_TILE_SIZE 0x10000u
+
 /* ID3D12Resource */
 struct d3d12_resource
 {
@@ -727,6 +747,10 @@ static inline bool d3d12_resource_is_texture(const struct d3d12_resource *resour

bool d3d12_resource_is_cpu_accessible(const struct d3d12_resource *resource);
 HRESULT d3d12_resource_validate_desc(const D3D12_RESOURCE_DESC *desc, struct d3d12_device *device);
+void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_resource *resource,
+        UINT *total_tile_count, D3D12_PACKED_MIP_INFO *packed_mip_info, D3D12_TILE_SHAPE *standard_tile_shape,
+        UINT *sub_resource_tiling_count, UINT first_sub_resource_tiling,
+        D3D12_SUBRESOURCE_TILING *sub_resource_tilings);

HRESULT d3d12_committed_resource_create(struct d3d12_device *device,
         const D3D12_HEAP_PROPERTIES *heap_properties, D3D12_HEAP_FLAGS heap_flags,
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 9f00d74fb..2801ec45a 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -36619,6 +36619,12 @@ static void test_get_resource_tiling(void)
             continue;
         }

+        if (tests[i].dim != D3D12_RESOURCE_DIMENSION_BUFFER)
+        {
+            skip("Tiled textures not supported.\n");
+            continue;
+        }
+
         memset(&packed_mip_info, 0xaa, sizeof(packed_mip_info));
         memset(&tile_shape, 0xaa, sizeof(tile_shape));
         memset(&tilings, 0xaa, sizeof(tilings));
@@ -36651,23 +36657,23 @@ static void test_get_resource_tiling(void)

ID3D12Device_GetResourceTiling(context.device, resource, &num_resource_tiles, &packed_mip_info, &tile_shape, &num_tilings, 0, tilings);

-        todo ok(num_resource_tiles != 0xdeadbeef && num_resource_tiles >= tests[i].expected_tile_count,
+        ok(num_resource_tiles != 0xdeadbeef && num_resource_tiles >= tests[i].expected_tile_count,
                 "Unexpected resource tile count %u.\n", num_resource_tiles);
-        todo ok(num_tilings == tests[i].expected_tiling_count, "Unexpected subresource tiling count %u.\n", num_tilings);
+        ok(num_tilings == tests[i].expected_tiling_count, "Unexpected subresource tiling count %u.\n", num_tilings);

-        todo ok(packed_mip_info.NumStandardMips != 0xaa && packed_mip_info.NumStandardMips >= tests[i].expected_standard_mips,
+        ok(packed_mip_info.NumStandardMips != 0xaa && packed_mip_info.NumStandardMips >= tests[i].expected_standard_mips,
                 "Unexpected standard mip count %u.\n", packed_mip_info.NumStandardMips);
-        todo ok(packed_mip_info.NumPackedMips == (tests[i].dim == D3D12_RESOURCE_DIMENSION_BUFFER
+        ok(packed_mip_info.NumPackedMips == (tests[i].dim == D3D12_RESOURCE_DIMENSION_BUFFER
                 ? 0 : tests[i].mip_levels - packed_mip_info.NumStandardMips),
                 "Unexpected packed mip count %u.\n", packed_mip_info.NumPackedMips);
-        todo ok(packed_mip_info.NumPackedMips != 0xaa
+        ok(packed_mip_info.NumPackedMips != 0xaa
                 && (packed_mip_info.NumTilesForPackedMips == 0) == (packed_mip_info.NumPackedMips == 0),
                 "Unexpected packed tile count %u.\n", packed_mip_info.NumTilesForPackedMips);

/* Docs say that tile shape should be cleared to zero if there is no standard mip, but drivers don't seem to care about this. */
-        todo ok(tile_shape.WidthInTexels == tests[i].tile_shape_w, "Unexpected tile width %u.\n", tile_shape.WidthInTexels);
-        todo ok(tile_shape.HeightInTexels == tests[i].tile_shape_h, "Unexpected tile height %u.\n", tile_shape.HeightInTexels);
-        todo ok(tile_shape.DepthInTexels == tests[i].tile_shape_d, "Unexpected tile depth %u.\n", tile_shape.DepthInTexels);
+        ok(tile_shape.WidthInTexels == tests[i].tile_shape_w, "Unexpected tile width %u.\n", tile_shape.WidthInTexels);
+        ok(tile_shape.HeightInTexels == tests[i].tile_shape_h, "Unexpected tile height %u.\n", tile_shape.HeightInTexels);
+        ok(tile_shape.DepthInTexels == tests[i].tile_shape_d, "Unexpected tile depth %u.\n", tile_shape.DepthInTexels);

for (j = 0; j < tests[i].expected_tiling_count; j++)
         {
@@ -36682,28 +36688,28 @@ static void test_get_resource_tiling(void)
                 if (tests[i].dim == D3D12_RESOURCE_DIMENSION_TEXTURE3D)
                     expected_d = compute_tile_count(tests[i].depth_or_array_layers, mip, tests[i].tile_shape_d);

-                todo ok(tilings[j].WidthInTiles == expected_w, "Unexpected width %u for subresource %u.\n", tilings[j].WidthInTiles, j);
-                todo ok(tilings[j].HeightInTiles == expected_h, "Unexpected width %u for subresource %u.\n", tilings[j].HeightInTiles, j);
-                todo ok(tilings[j].DepthInTiles == expected_d, "Unexpected width %u for subresource %u.\n", tilings[j].DepthInTiles, j);
+                ok(tilings[j].WidthInTiles == expected_w, "Unexpected width %u for subresource %u.\n", tilings[j].WidthInTiles, j);
+                ok(tilings[j].HeightInTiles == expected_h, "Unexpected width %u for subresource %u.\n", tilings[j].HeightInTiles, j);
+                ok(tilings[j].DepthInTiles == expected_d, "Unexpected width %u for subresource %u.\n", tilings[j].DepthInTiles, j);

-                todo ok(tilings[j].StartTileIndexInOverallResource == tile_index, "Unexpected start tile index %u for subresource %u.\n",
+                ok(tilings[j].StartTileIndexInOverallResource == tile_index, "Unexpected start tile index %u for subresource %u.\n",
                         tilings[j].StartTileIndexInOverallResource, j);

tile_index += tilings[j].WidthInTiles * tilings[j].HeightInTiles * tilings[j].DepthInTiles;
             }
             else
             {
-                todo ok(!tilings[j].WidthInTiles && !tilings[j].HeightInTiles && !tilings[j].DepthInTiles,
+                ok(!tilings[j].WidthInTiles && !tilings[j].HeightInTiles && !tilings[j].DepthInTiles,
                         "Unexpected tile count (%u,%u,%u) for packed subresource %u.\n",
                         tilings[j].WidthInTiles, tilings[j].HeightInTiles, tilings[j].DepthInTiles, j);
-                todo ok(tilings[j].StartTileIndexInOverallResource == 0xffffffff, "Unexpected start tile index %u for packed subresource %u.\n",
+                ok(tilings[j].StartTileIndexInOverallResource == 0xffffffff, "Unexpected start tile index %u for packed subresource %u.\n",
                         tilings[j].StartTileIndexInOverallResource, j);
             }
         }

-        todo ok(num_resource_tiles == tile_index + packed_mip_info.NumTilesForPackedMips,
+        ok(num_resource_tiles == tile_index + packed_mip_info.NumTilesForPackedMips,
                 "Unexpected resource tile count %u.\n", num_resource_tiles);
-        todo ok(packed_mip_info.StartTileIndexInOverallResource == (packed_mip_info.NumPackedMips ? tile_index : 0),
+        ok(packed_mip_info.StartTileIndexInOverallResource == (packed_mip_info.NumPackedMips ? tile_index : 0),
                 "Unexpected mip tail start tile index %u.\n", packed_mip_info.StartTileIndexInOverallResource);

ID3D12Resource_Release(resource);

[PATCH 10/21] vkd3d: Implement ID3D12Device::GetResourceTiling() for textures.

[Conor McCarthy]

From: Conor McCarthy

---
 include/vkd3d_d3d12.idl    |  1 +
 libs/vkd3d/device.c        |  6 ---
 libs/vkd3d/resource.c      | 90 ++++++++++++++++++++++++++++++++++----
 libs/vkd3d/vkd3d_private.h |  1 +
 tests/d3d12.c              | 14 ++----
 5 files changed, 87 insertions(+), 25 deletions(-)

diff --git a/include/vkd3d_d3d12.idl b/include/vkd3d_d3d12.idl
index 06287f5a8..495a04678 100644
--- a/include/vkd3d_d3d12.idl
+++ b/include/vkd3d_d3d12.idl
@@ -75,6 +75,7 @@ const UINT D3D12_TEXTURE_DATA_PLACEMENT_ALIGNMENT = 512;
 const UINT D3D12_UAV_COUNTER_PLACEMENT_ALIGNMENT = 4096;
 const UINT D3D12_VS_INPUT_REGISTER_COUNT = 32;
 const UINT D3D12_VIEWPORT_AND_SCISSORRECT_OBJECT_COUNT_PER_PIPELINE = 16;
+const UINT D3D12_PACKED_TILE = 0xffffffff;

cpp_quote("#endif")

diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 69dbb8705..b1eba634f 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -3900,12 +3900,6 @@ static void STDMETHODCALLTYPE d3d12_device_GetResourceTiling(ID3D12Device *iface
             sub_resource_tiling_count, first_sub_resource_tiling,
             sub_resource_tilings);

-    if (d3d12_resource_is_texture(resource_impl))
-    {
-        FIXME("Not implemented for textures.\n");
-        return;
-    }
-
     d3d12_resource_get_tiling(device, resource_impl, total_tile_count, packed_mip_info, standard_tile_shape,
             sub_resource_tiling_count, first_sub_resource_tiling, sub_resource_tilings);
 }
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index c50d6241f..a6e76f197 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1065,12 +1065,26 @@ static void d3d12_resource_get_level_box(const struct d3d12_resource *resource,
     box->back = d3d12_resource_desc_get_depth(&resource->desc, level);
 }

+static void compute_image_subresource_size_in_tiles(const VkExtent3D *tile_extent,
+        const struct D3D12_RESOURCE_DESC *desc, unsigned int miplevel_idx,
+        struct vkd3d_tiled_region_extent *size)
+{
+    unsigned int width, height, depth;
+
+    width = d3d12_resource_desc_get_width(desc, miplevel_idx);
+    height = d3d12_resource_desc_get_height(desc, miplevel_idx);
+    depth = d3d12_resource_desc_get_depth(desc, miplevel_idx);
+    size->width = (width + tile_extent->width - 1) / tile_extent->width;
+    size->height = (height + tile_extent->height - 1) / tile_extent->height;
+    size->depth = (depth + tile_extent->depth - 1) / tile_extent->depth;
+}
+
 void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_resource *resource,
         UINT *total_tile_count, D3D12_PACKED_MIP_INFO *packed_mip_info, D3D12_TILE_SHAPE *standard_tile_shape,
         UINT *subresource_tiling_count, UINT first_subresource_tiling,
         D3D12_SUBRESOURCE_TILING *subresource_tilings)
 {
-    unsigned int i, subresource, subresource_count, count;
+    unsigned int i, subresource, subresource_count, miplevel_idx, count;
     const struct vkd3d_subresource_tile_info *tile_info;
     const VkExtent3D *tile_extent;

@@ -1079,9 +1093,10 @@ void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_r
     if (packed_mip_info)
     {
         packed_mip_info->NumStandardMips = resource->tiles.standard_mip_count;
-        packed_mip_info->NumPackedMips = 0;
-        packed_mip_info->NumTilesForPackedMips = 0;
-        packed_mip_info->StartTileIndexInOverallResource = 0;
+        packed_mip_info->NumPackedMips = resource->desc.MipLevels - packed_mip_info->NumStandardMips;
+        packed_mip_info->NumTilesForPackedMips = !!resource->tiles.packed_mip_tile_count; /* non-zero dummy value */
+        packed_mip_info->StartTileIndexInOverallResource = packed_mip_info->NumPackedMips
+                ? resource->tiles.subresources[resource->tiles.standard_mip_count].offset : 0;
     }

if (standard_tile_shape)
@@ -1106,6 +1121,14 @@ void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_r
     for (i = 0; i < count; ++i)
     {
         subresource = i + first_subresource_tiling;
+        miplevel_idx = subresource % resource->desc.MipLevels;
+        if (miplevel_idx >= resource->tiles.standard_mip_count)
+        {
+            memset(&subresource_tilings[i], 0, sizeof(subresource_tilings[i]));
+            subresource_tilings[i].StartTileIndexInOverallResource = D3D12_PACKED_TILE;
+            continue;
+        }
+
         tile_info = &resource->tiles.subresources[subresource];
         subresource_tilings[i].StartTileIndexInOverallResource = tile_info->offset;
         subresource_tilings[i].WidthInTiles = tile_info->extent.width;
@@ -1117,14 +1140,17 @@ void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_r

static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3d12_device *device)
 {
+    unsigned int i, start_idx, subresource_count, tile_count, miplevel_idx;
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkSparseImageMemoryRequirements sparse_requirements_buf[3];
+    VkSparseImageMemoryRequirements sparse_requirements;
     struct vkd3d_subresource_tile_info *tile_info;
-    unsigned int subresource_count;
+    VkMemoryRequirements requirements;
+    const VkExtent3D *tile_extent;
+    uint32_t requirement_count;

subresource_count = d3d12_resource_desc_get_sub_resource_count(&resource->desc);

-    if (!d3d12_resource_is_buffer(resource))
-        return true;
-
     if (!(resource->tiles.subresources = vkd3d_calloc(subresource_count, sizeof(*resource->tiles.subresources))))
     {
         ERR("Failed to allocate subresource info array.\n");
@@ -1146,10 +1172,56 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
         resource->tiles.total_count = tile_info->extent.width;
         resource->tiles.subresource_count = 1;
         resource->tiles.standard_mip_count = 1;
+        resource->tiles.packed_mip_tile_count = 0;
     }
     else
     {
-        vkd3d_unreachable();
+        VK_CALL(vkGetImageMemoryRequirements(device->vk_device, resource->u.vk_image, &requirements));
+        if (requirements.alignment > D3D12_TILE_SIZE)
+            FIXME("Vulkan device tile size is greater than the standard D3D12 tile size.\n");
+
+        requirement_count = ARRAY_SIZE(sparse_requirements_buf);
+        VK_CALL(vkGetImageSparseMemoryRequirements(device->vk_device, resource->u.vk_image,
+                &requirement_count, sparse_requirements_buf));
+        for (i = 0; i < requirement_count; ++i)
+        {
+            if (!(sparse_requirements_buf[i].formatProperties.aspectMask & VK_IMAGE_ASPECT_METADATA_BIT))
+            {
+                sparse_requirements = sparse_requirements_buf[i];
+                break;
+            }
+        }
+
+        resource->tiles.tile_extent = sparse_requirements.formatProperties.imageGranularity;
+        resource->tiles.subresource_count = subresource_count;
+        resource->tiles.standard_mip_count = sparse_requirements.imageMipTailSize
+                ? sparse_requirements.imageMipTailFirstLod : resource->desc.MipLevels;
+        resource->tiles.packed_mip_tile_count = (resource->tiles.standard_mip_count < resource->desc.MipLevels)
+                ? sparse_requirements.imageMipTailSize / requirements.alignment : 0;
+
+        for (i = 0, start_idx = 0; i < subresource_count; ++i)
+        {
+            miplevel_idx = i % resource->desc.MipLevels;
+
+            tile_extent = &sparse_requirements.formatProperties.imageGranularity;
+            tile_info = &resource->tiles.subresources[i];
+            compute_image_subresource_size_in_tiles(tile_extent, &resource->desc, miplevel_idx, &tile_info->extent);
+            tile_info->offset = start_idx;
+            tile_info->count = 0;
+
+            if (miplevel_idx < resource->tiles.standard_mip_count)
+            {
+                tile_count = tile_info->extent.width * tile_info->extent.height * tile_info->extent.depth;
+                start_idx += tile_count;
+                tile_info->count = tile_count;
+            }
+            else if (miplevel_idx == resource->tiles.standard_mip_count)
+            {
+                tile_info->count = 1; /* Non-zero dummy value */
+                start_idx += 1;
+            }
+        }
+        resource->tiles.total_count = start_idx;
     }

return true;
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index b545945a5..859844684 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -691,6 +691,7 @@ struct d3d12_resource_tile_info
     VkExtent3D tile_extent;
     unsigned int total_count;
     unsigned int standard_mip_count;
+    unsigned int packed_mip_tile_count;
     unsigned int subresource_count;
     struct vkd3d_subresource_tile_info *subresources;
 };
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 2801ec45a..6db1a79d0 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -36588,18 +36588,18 @@ static void test_get_resource_tiling(void)

num_tilings = ARRAY_SIZE(tilings);
     ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 10, tilings);
-    todo ok(num_tilings == 0, "Unexpected tiling count %u.\n", num_tilings);
+    ok(num_tilings == 0, "Unexpected tiling count %u.\n", num_tilings);

num_tilings = ARRAY_SIZE(tilings);
     ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 2, tilings);
-    todo ok(num_tilings == 8, "Unexpected tiling count %u.\n", num_tilings);
-    todo ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);
+    ok(num_tilings == 8, "Unexpected tiling count %u.\n", num_tilings);
+    ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);

num_tilings = 1;
     memset(&tilings, 0xaa, sizeof(tilings));
     ID3D12Device_GetResourceTiling(context.device, resource, NULL, NULL, NULL, &num_tilings, 2, tilings);
     ok(num_tilings == 1, "Unexpected tiling count %u.\n", num_tilings);
-    todo ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);
+    ok(tilings[0].StartTileIndexInOverallResource == 20, "Unexpected start tile index %u.\n", tilings[0].StartTileIndexInOverallResource);
     ok(tilings[1].StartTileIndexInOverallResource == 0xaaaaaaaa, "Tiling array got modified.\n");

ID3D12Resource_Release(resource);
@@ -36619,12 +36619,6 @@ static void test_get_resource_tiling(void)
             continue;
         }

-        if (tests[i].dim != D3D12_RESOURCE_DIMENSION_BUFFER)
-        {
-            skip("Tiled textures not supported.\n");
-            continue;
-        }
-
         memset(&packed_mip_info, 0xaa, sizeof(packed_mip_info));
         memset(&tile_shape, 0xaa, sizeof(tile_shape));
         memset(&tilings, 0xaa, sizeof(tilings));

[PATCH 11/21] tests/d3d12: Add tests for UpdateTileMappings().

[Conor McCarthy]

From: Conor McCarthy

Based on a vkd3d-proton patch by Philip Rebohle.
---
 tests/d3d12.c | 830 ++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 830 insertions(+)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 6db1a79d0..083f9c4d5 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -36712,6 +36712,835 @@ static void test_get_resource_tiling(void)
     destroy_test_context(&context);
 }

+
+static void set_region_offset(D3D12_TILED_RESOURCE_COORDINATE *region, uint32_t x, uint32_t y, uint32_t z, uint32_t subresource)
+{
+    region->X = x;
+    region->Y = y;
+    region->Z = z;
+    region->Subresource = subresource;
+}
+
+static void set_region_size(D3D12_TILE_REGION_SIZE *region, uint32_t num_tiles, bool use_box, uint32_t w, uint32_t h, uint32_t d)
+{
+    region->NumTiles = num_tiles;
+    region->UseBox = use_box;
+    region->Width = w;
+    region->Height = h;
+    region->Depth = d;
+}
+
+static void test_update_tile_mappings(void)
+{
+    ID3D12Resource *resource, *resource_2, *array_resource, *readback_buffer;
+    D3D12_RESOURCE_DESC resource_desc, array_resource_desc;
+    D3D12_TILED_RESOURCE_COORDINATE region_offsets[8];
+    ID3D12PipelineState *check_texture_3d_pipeline;
+    D3D12_ROOT_SIGNATURE_DESC root_signature_desc;
+    ID3D12PipelineState *clear_texture_pipeline;
+    ID3D12PipelineState *check_texture_pipeline;
+    ID3D12PipelineState *check_buffer_pipeline;
+    D3D12_UNORDERED_ACCESS_VIEW_DESC uav_desc;
+    ID3D12DescriptorHeap *cpu_heap, *gpu_heap;
+    ID3D12RootSignature *clear_root_signature;
+    D3D12_SHADER_RESOURCE_VIEW_DESC srv_desc;
+    D3D12_DESCRIPTOR_RANGE descriptor_range;
+    D3D12_ROOT_PARAMETER root_parameters[2];
+    D3D12_TILE_REGION_SIZE region_sizes[8];
+    D3D12_GPU_VIRTUAL_ADDRESS readback_va;
+    D3D12_HEAP_PROPERTIES heap_properties;
+    D3D12_PACKED_MIP_INFO packed_mip_info;
+    D3D12_SUBRESOURCE_TILING tilings[10];
+    D3D12_TILE_RANGE_FLAGS tile_flags[8];
+    ID3D12RootSignature *root_signature;
+    struct d3d12_resource_readback rb;
+    struct test_context_desc desc;
+    struct test_context context;
+    D3D12_TILE_SHAPE tile_shape;
+    unsigned int i, j, x, y, z;
+    D3D12_HEAP_DESC heap_desc;
+    UINT tile_offsets[8];
+    UINT tile_counts[8];
+    ID3D12Heap *heap;
+    UINT num_tilings;
+    D3D12_BOX box;
+    HRESULT hr;
+
+#if 0
+    StructuredBuffer<uint> tiled_buffer : register(t0);
+    RWStructuredBuffer<uint> out_buffer : register(u0);
+
+    [numthreads(64, 1, 1)]
+    void main(uint3 thread_id : SV_DispatchThreadID)
+    {
+        out_buffer[thread_id.x] = tiled_buffer[16384 * thread_id.x];
+    }
+#endif
+    static const DWORD cs_buffer_code[] =
+    {
+        0x43425844, 0xa8625c41, 0xfd85df89, 0xcedb7945, 0x0e3444ea, 0x00000001, 0x00000108, 0x00000003,
+        0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+        0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000000b4, 0x00050050, 0x0000002d, 0x0100086a,
+        0x040000a2, 0x00107000, 0x00000000, 0x00000004, 0x0400009e, 0x0011e000, 0x00000000, 0x00000004,
+        0x0200005f, 0x00020012, 0x02000068, 0x00000001, 0x0400009b, 0x00000040, 0x00000001, 0x00000001,
+        0x06000029, 0x00100012, 0x00000000, 0x0002000a, 0x00004001, 0x0000000e, 0x8b0000a7, 0x80002302,
+        0x00199983, 0x00100012, 0x00000000, 0x0010000a, 0x00000000, 0x00004001, 0x00000000, 0x00107006,
+        0x00000000, 0x080000a8, 0x0011e012, 0x00000000, 0x0002000a, 0x00004001, 0x00000000, 0x0010000a,
+        0x00000000, 0x0100003e,
+    };
+#if 0
+    Texture2D<uint> tiled_texture : register(t0);
+    RWStructuredBuffer<uint> out_buffer : register(u0);
+
+    [numthreads(28,1,1)]
+    void main(uint3 thread_id : SV_DispatchThreadID)
+    {
+        uint2 tile_size = uint2(128, 128);
+        uint tile_index = 0;
+        uint tile_count = 4;
+        uint mip_count = 10;
+        uint mip_level = 0;
+
+        while (thread_id.x >= tile_index + tile_count * tile_count && mip_level < mip_count)
+        {
+            tile_index += tile_count * tile_count;
+            tile_count = max(tile_count / 2, 1);
+            mip_level += 1;
+        }
+
+        uint2 tile_coord;
+        tile_coord.x = (thread_id.x - tile_index) % tile_count;
+        tile_coord.y = (thread_id.x - tile_index) / tile_count;
+
+        out_buffer[thread_id.x] = tiled_texture.mips[mip_level][tile_coord * tile_size];
+    }
+#endif
+    static const DWORD cs_texture_code[] =
+    {
+        0x43425844, 0x03e118db, 0xda7deb90, 0xedb39031, 0x6b646a0b, 0x00000001, 0x00000288, 0x00000003,
+        0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+        0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x00000234, 0x00050050, 0x0000008d, 0x0100086a,
+        0x04001858, 0x00107000, 0x00000000, 0x00004444, 0x0400009e, 0x0011e000, 0x00000000, 0x00000004,
+        0x0200005f, 0x00020012, 0x02000068, 0x00000003, 0x0400009b, 0x0000001c, 0x00000001, 0x00000001,
+        0x08000036, 0x00100072, 0x00000000, 0x00004002, 0x00000000, 0x00000004, 0x00000000, 0x00000000,
+        0x01000030, 0x09000023, 0x00100082, 0x00000000, 0x0010001a, 0x00000000, 0x0010001a, 0x00000000,
+        0x0010000a, 0x00000000, 0x06000050, 0x00100012, 0x00000001, 0x0002000a, 0x0010003a, 0x00000000,
+        0x0700004f, 0x00100022, 0x00000001, 0x0010002a, 0x00000000, 0x00004001, 0x0000000a, 0x07000001,
+        0x00100012, 0x00000001, 0x0010001a, 0x00000001, 0x0010000a, 0x00000001, 0x03000003, 0x0010000a,
+        0x00000001, 0x07000055, 0x00100012, 0x00000001, 0x0010001a, 0x00000000, 0x00004001, 0x00000001,
+        0x07000053, 0x00100022, 0x00000000, 0x0010000a, 0x00000001, 0x00004001, 0x00000001, 0x0700001e,
+        0x00100042, 0x00000000, 0x0010002a, 0x00000000, 0x00004001, 0x00000001, 0x05000036, 0x00100012,
+        0x00000000, 0x0010003a, 0x00000000, 0x01000016, 0x05000036, 0x001000c2, 0x00000001, 0x00100aa6,
+        0x00000000, 0x0700001e, 0x00100012, 0x00000000, 0x8010000a, 0x00000041, 0x00000000, 0x0002000a,
+        0x0900004e, 0x00100012, 0x00000000, 0x00100012, 0x00000002, 0x0010000a, 0x00000000, 0x0010001a,
+        0x00000000, 0x05000036, 0x00100022, 0x00000002, 0x0010000a, 0x00000000, 0x0a000029, 0x00100032,
+        0x00000001, 0x00100046, 0x00000002, 0x00004002, 0x00000007, 0x00000007, 0x00000000, 0x00000000,
+        0x8900002d, 0x800000c2, 0x00111103, 0x00100012, 0x00000000, 0x00100e46, 0x00000001, 0x00107e46,
+        0x00000000, 0x080000a8, 0x0011e012, 0x00000000, 0x0002000a, 0x00004001, 0x00000000, 0x0010000a,
+        0x00000000, 0x0100003e,
+    };
+
+#if 0
+    Texture3D<uint> tiled_texture : register(t0);
+    RWStructuredBuffer<uint> out_buffer : register(u0);
+
+    [numthreads(9,1,1)]
+    void main(uint3 thread_id : SV_DispatchThreadID)
+    {
+        uint3 tile_size = uint3(32, 32, 16);
+        uint tile_index = 0;
+        uint tile_count = 2;
+        uint mip_count = 2;
+        uint mip_level = 0;
+
+        while (thread_id.x >= tile_index + tile_count * tile_count * tile_count && mip_level < mip_count)
+        {
+            tile_index += tile_count * tile_count * tile_count;
+            tile_count = max(tile_count / 2, 1);
+            mip_level += 1;
+        }
+
+        uint3 tile_coord;
+        tile_coord.x = (thread_id.x - tile_index) % tile_count;
+        tile_coord.y = ((thread_id.x - tile_index) / tile_count) % tile_count;
+        tile_coord.z = (thread_id.x - tile_index) / (tile_count * tile_count);
+
+        out_buffer[thread_id.x] = tiled_texture.mips[mip_level][tile_coord * tile_size];
+    }
+#endif
+    static const DWORD cs_texture_3d_code[] =
+    {
+        0x43425844, 0x71b4eb36, 0x2c65e68d, 0x7763693f, 0xfd4eafc6, 0x00000001, 0x000002f4, 0x00000003,
+        0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+        0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x000002a0, 0x00050050, 0x000000a8, 0x0100086a,
+        0x04002858, 0x00107000, 0x00000000, 0x00004444, 0x0400009e, 0x0011e000, 0x00000000, 0x00000004,
+        0x0200005f, 0x00020012, 0x02000068, 0x00000004, 0x0400009b, 0x00000009, 0x00000001, 0x00000001,
+        0x08000036, 0x00100032, 0x00000000, 0x00004002, 0x00000000, 0x00000002, 0x00000000, 0x00000000,
+        0x05000036, 0x00100082, 0x00000001, 0x00004001, 0x00000000, 0x01000030, 0x08000026, 0x0000d000,
+        0x00100042, 0x00000000, 0x0010001a, 0x00000000, 0x0010001a, 0x00000000, 0x09000023, 0x00100042,
+        0x00000000, 0x0010002a, 0x00000000, 0x0010001a, 0x00000000, 0x0010000a, 0x00000000, 0x06000050,
+        0x00100082, 0x00000000, 0x0002000a, 0x0010002a, 0x00000000, 0x0700004f, 0x00100012, 0x00000002,
+        0x0010003a, 0x00000001, 0x00004001, 0x00000002, 0x07000001, 0x00100082, 0x00000000, 0x0010003a,
+        0x00000000, 0x0010000a, 0x00000002, 0x03000003, 0x0010003a, 0x00000000, 0x07000055, 0x00100082,
+        0x00000000, 0x0010001a, 0x00000000, 0x00004001, 0x00000001, 0x07000053, 0x00100022, 0x00000000,
+        0x0010003a, 0x00000000, 0x00004001, 0x00000001, 0x0700001e, 0x00100082, 0x00000001, 0x0010003a,
+        0x00000001, 0x00004001, 0x00000001, 0x05000036, 0x00100012, 0x00000000, 0x0010002a, 0x00000000,
+        0x01000016, 0x0700001e, 0x00100012, 0x00000000, 0x8010000a, 0x00000041, 0x00000000, 0x0002000a,
+        0x0900004e, 0x00100012, 0x00000002, 0x00100012, 0x00000003, 0x0010000a, 0x00000000, 0x0010001a,
+        0x00000000, 0x0800004e, 0x0000d000, 0x00100022, 0x00000003, 0x0010000a, 0x00000002, 0x0010001a,
+        0x00000000, 0x08000026, 0x0000d000, 0x00100022, 0x00000000, 0x0010001a, 0x00000000, 0x0010001a,
+        0x00000000, 0x0800004e, 0x00100042, 0x00000003, 0x0000d000, 0x0010000a, 0x00000000, 0x0010001a,
+        0x00000000, 0x0a000029, 0x00100072, 0x00000001, 0x00100246, 0x00000003, 0x00004002, 0x00000005,
+        0x00000005, 0x00000004, 0x00000000, 0x8900002d, 0x80000142, 0x00111103, 0x00100012, 0x00000000,
+        0x00100e46, 0x00000001, 0x00107e46, 0x00000000, 0x080000a8, 0x0011e012, 0x00000000, 0x0002000a,
+        0x00004001, 0x00000000, 0x0010000a, 0x00000000, 0x0100003e,
+    };
+
+#if 0
+    RWTexture3D<uint> uav : register(u0);
+
+    cbuffer clear_args
+    {
+        uint3 offset;
+        uint value;
+    };
+
+    [numthreads(4, 4, 4)]
+    void main(uint3 coord : SV_DispatchThreadID)
+    {
+        uav[offset + coord] = value;
+    }
+#endif
+    static const DWORD cs_clear_code[] =
+    {
+        0x43425844, 0x288d0bcd, 0xbe5e644d, 0x95665c2e, 0xd8f02c36, 0x00000001, 0x000000e0, 0x00000003,
+        0x0000002c, 0x0000003c, 0x0000004c, 0x4e475349, 0x00000008, 0x00000000, 0x00000008, 0x4e47534f,
+        0x00000008, 0x00000000, 0x00000008, 0x58454853, 0x0000008c, 0x00050050, 0x00000023, 0x0100086a,
+        0x04000059, 0x00208e46, 0x00000000, 0x00000001, 0x0400289c, 0x0011e000, 0x00000000, 0x00004444,
+        0x0200005f, 0x00020072, 0x02000068, 0x00000001, 0x0400009b, 0x00000004, 0x00000004, 0x00000004,
+        0x0700001e, 0x001000f2, 0x00000000, 0x00020a46, 0x00208a46, 0x00000000, 0x00000000, 0x080000a4,
+        0x0011e0f2, 0x00000000, 0x00100e46, 0x00000000, 0x00208ff6, 0x00000000, 0x00000000, 0x0100003e,
+    };
+
+    static const D3D12_SHADER_BYTECODE cs_texture = { cs_texture_code, sizeof(cs_texture_code) };
+    static const D3D12_SHADER_BYTECODE cs_texture_3d = { cs_texture_3d_code, sizeof(cs_texture_3d_code) };
+    static const D3D12_SHADER_BYTECODE cs_buffer = { cs_buffer_code, sizeof(cs_buffer_code) };
+    static const D3D12_SHADER_BYTECODE cs_clear = { cs_clear_code, sizeof(cs_clear_code) };
+
+    static const uint32_t buffer_region_tiles[] =
+    {
+    /*     0   1   2   3   4   5   6   7   8   9 */
+    /*0*/ 33, 34, 35, 36, 37,  6,  7,  8,  9, 10,
+    /*1*/ 11, 12, 38, 39, 40, 41,  1, 18,  2, 20,
+    /*2*/ 21, 22, 23,  3,  4,  4,  4,  0,  0, 25,
+    /*3*/ 26, 27, 28, 29, 30, 36, 37, 38, 39, 40,
+    /*4*/  9, 11, 43, 44, 45, 46, 45, 46, 49, 50,
+    /*5*/  0,  0, 17, 18, 19, 20, 21, 22, 23, 24,
+    /*6*/ 61, 62, 63, 12,
+    };
+
+    static const uint32_t texture_region_tiles[] =
+    {
+        1, 2, 4, 5, 6, 7, 1, 1, 9, 1, 17, 14, 8, 14, 3, 0,
+        18, 18, 19, 18, 19, 22, 23, 24, 25, 26, 27, 28,
+    };
+
+    static const uint32_t texture_3d_region_tiles[] =
+    {
+        3, 2, 0, 7, 8, 2, 4, 5, 6,
+    };
+
+    memset(&desc, 0, sizeof(desc));
+    desc.rt_width = 640;
+    desc.rt_height = 480;
+    desc.rt_format = DXGI_FORMAT_R8G8B8A8_UNORM;
+    if (!init_test_context(&context, &desc))
+        return;
+
+    descriptor_range.RangeType = D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
+    descriptor_range.NumDescriptors = 1;
+    descriptor_range.BaseShaderRegister = 0;
+    descriptor_range.RegisterSpace = 0;
+    descriptor_range.OffsetInDescriptorsFromTableStart = 0;
+    root_parameters[0].ParameterType = D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE;
+    root_parameters[0].DescriptorTable.NumDescriptorRanges = 1;
+    root_parameters[0].DescriptorTable.pDescriptorRanges = &descriptor_range;
+    root_parameters[0].ShaderVisibility = D3D12_SHADER_VISIBILITY_ALL;
+    root_parameters[1].ParameterType = D3D12_ROOT_PARAMETER_TYPE_UAV;
+    root_parameters[1].Descriptor.ShaderRegister = 0;
+    root_parameters[1].Descriptor.RegisterSpace = 0;
+    root_parameters[1].ShaderVisibility = D3D12_SHADER_VISIBILITY_ALL;
+    root_signature_desc.NumParameters = ARRAY_SIZE(root_parameters);
+    root_signature_desc.pParameters = root_parameters;
+    root_signature_desc.NumStaticSamplers = 0;
+    root_signature_desc.pStaticSamplers = NULL;
+    root_signature_desc.Flags = D3D12_ROOT_SIGNATURE_FLAG_NONE;
+    hr = create_root_signature(context.device, &root_signature_desc, &root_signature);
+    ok(hr == S_OK, "Failed to create root signature, hr %#x.\n", hr);
+
+    descriptor_range.RangeType = D3D12_DESCRIPTOR_RANGE_TYPE_UAV;
+    root_parameters[1].ParameterType = D3D12_ROOT_PARAMETER_TYPE_32BIT_CONSTANTS;
+    root_parameters[1].Constants.ShaderRegister = 0;
+    root_parameters[1].Constants.RegisterSpace = 0;
+    root_parameters[1].Constants.Num32BitValues = 4;
+    root_parameters[1].ShaderVisibility = D3D12_SHADER_VISIBILITY_ALL;
+    hr = create_root_signature(context.device, &root_signature_desc, &clear_root_signature);
+    ok(hr == S_OK, "Failed to create root signature, hr %#x.\n", hr);
+
+    clear_texture_pipeline = create_compute_pipeline_state(context.device, clear_root_signature, cs_clear);
+    check_texture_pipeline = create_compute_pipeline_state(context.device, root_signature, cs_texture);
+    check_texture_3d_pipeline = create_compute_pipeline_state(context.device, root_signature, cs_texture_3d);
+    check_buffer_pipeline = create_compute_pipeline_state(context.device, root_signature, cs_buffer);
+
+    cpu_heap = create_cpu_descriptor_heap(context.device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 11);
+    gpu_heap = create_gpu_descriptor_heap(context.device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 11);
+
+    memset(&heap_properties, 0, sizeof(heap_properties));
+    heap_properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
+    resource_desc.Alignment = 0;
+    resource_desc.Width = 64 * sizeof(uint32_t);
+    resource_desc.Height = 1;
+    resource_desc.DepthOrArraySize = 1;
+    resource_desc.MipLevels = 1;
+    resource_desc.Format = DXGI_FORMAT_UNKNOWN;
+    resource_desc.SampleDesc.Count = 1;
+    resource_desc.SampleDesc.Quality = 0;
+    resource_desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+    resource_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+    hr = ID3D12Device_CreateCommittedResource(context.device, &heap_properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&readback_buffer);
+    ok(hr == S_OK, "Failed to create readback buffer, hr %#x.\n", hr);
+
+    readback_va = ID3D12Resource_GetGPUVirtualAddress(readback_buffer);
+
+    /* Test buffer tile mappings */
+    heap_desc.Properties = heap_properties;
+    heap_desc.Alignment = 0;
+    heap_desc.SizeInBytes = 64 * 65536;
+    heap_desc.Flags = D3D12_HEAP_FLAG_ALLOW_ONLY_BUFFERS;
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void **)&heap);
+    ok(hr == S_OK, "Failed to create heap, hr %#x.\n", hr);
+
+    resource_desc.Width = 64 * 65536;
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == S_OK, "Failed to create reserved buffer, hr %#x.\n", hr);
+
+    srv_desc.Format = DXGI_FORMAT_UNKNOWN;
+    srv_desc.ViewDimension = D3D12_SRV_DIMENSION_BUFFER;
+    srv_desc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
+    srv_desc.Buffer.FirstElement = 0;
+    srv_desc.Buffer.NumElements = resource_desc.Width / sizeof(uint32_t);
+    srv_desc.Buffer.StructureByteStride = sizeof(uint32_t);
+    srv_desc.Buffer.Flags = D3D12_BUFFER_SRV_FLAG_NONE;
+    ID3D12Device_CreateShaderResourceView(context.device, resource, &srv_desc, get_cpu_descriptor_handle(&context, gpu_heap, 0));
+
+    uav_desc.Format = DXGI_FORMAT_R32_UINT;
+    uav_desc.ViewDimension = D3D12_UAV_DIMENSION_BUFFER;
+    uav_desc.Buffer.FirstElement = 0;
+    uav_desc.Buffer.NumElements = resource_desc.Width / sizeof(uint32_t);
+    uav_desc.Buffer.StructureByteStride = 0;
+    uav_desc.Buffer.CounterOffsetInBytes = 0;
+    uav_desc.Buffer.Flags = D3D12_BUFFER_UAV_FLAG_NONE;
+    ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, cpu_heap, 1));
+    ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, gpu_heap, 1));
+
+    /* Map entire buffer, linearly, and initialize tile data */
+    tile_offsets[0] = 0;
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource, 1, NULL, NULL,
+            heap, 1, NULL, tile_offsets, NULL, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    for (i = 0; i < 64; i++)
+    {
+        UINT clear_value[4] = { 0, 0, 0, 0 };
+        D3D12_RECT clear_rect;
+
+        set_rect(&clear_rect, 16384 * i, 0, 16384 * (i + 1), 1);
+        clear_value[0] = i + 1;
+
+        ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(context.list,
+                get_gpu_descriptor_handle(&context, gpu_heap, 1),
+                get_cpu_descriptor_handle(&context, cpu_heap, 1),
+                resource, clear_value, 1, &clear_rect);
+    }
+
+    transition_resource_state(context.list, resource, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE);
+    ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+    ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+    ID3D12GraphicsCommandList_SetPipelineState(context.list, check_buffer_pipeline);
+    ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+    ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+    ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < 64; i++)
+    {
+        set_box(&box, i, 0, 0, i + 1, 1, 1);
+        todo check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+    }
+
+    release_resource_readback(&rb);
+
+    /* Test arbitrary tile mappings */
+    set_region_offset(&region_offsets[0], 16, 0, 0, 0);
+    set_region_offset(&region_offsets[1], 18, 0, 0, 0);
+    set_region_offset(&region_offsets[2], 23, 0, 0, 0);
+    set_region_offset(&region_offsets[3], 40, 0, 0, 0);
+    set_region_offset(&region_offsets[4], 41, 0, 0, 0);
+    set_region_offset(&region_offsets[5], 63, 0, 0, 0);
+
+    tile_offsets[0] = 0;
+    tile_offsets[1] = 8;
+    tile_offsets[2] = 10;
+
+    tile_counts[0] = 3;
+    tile_counts[1] = 1;
+    tile_counts[2] = 2;
+
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource, 6, region_offsets, NULL,
+            heap, 3, NULL, tile_offsets, tile_counts, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[0], 24, 0, 0, 0);
+    set_region_offset(&region_offsets[1], 50, 0, 0, 0);
+    set_region_offset(&region_offsets[2], 0, 0, 0, 0);
+    set_region_offset(&region_offsets[3], 52, 0, 0, 0);
+    set_region_offset(&region_offsets[4], 29, 0, 0, 0);
+
+    set_region_size(&region_sizes[0], 5, false, 0, 0, 0);
+    set_region_size(&region_sizes[1], 2, false, 0, 0, 0);
+    set_region_size(&region_sizes[2], 16, false, 0, 0, 0);
+    set_region_size(&region_sizes[3], 8, false, 0, 0, 0);
+    set_region_size(&region_sizes[4], 6, false, 0, 0, 0);
+
+    tile_flags[0] = D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE;
+    tile_flags[1] = D3D12_TILE_RANGE_FLAG_NULL;
+    tile_flags[2] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[3] = D3D12_TILE_RANGE_FLAG_SKIP;
+    tile_flags[4] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[5] = D3D12_TILE_RANGE_FLAG_NONE;
+
+    tile_offsets[0] = 3;
+    tile_offsets[1] = 0;
+    tile_offsets[2] = 32;
+    tile_offsets[3] = 0;
+    tile_offsets[4] = 37;
+    tile_offsets[5] = 16;
+
+    tile_counts[0] = 3;
+    tile_counts[1] = 4;
+    tile_counts[2] = 5;
+    tile_counts[3] = 7;
+    tile_counts[4] = 4;
+    tile_counts[5] = 14;
+
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource, 5, region_offsets, region_sizes,
+            heap, 6, tile_flags, tile_offsets, tile_counts, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[0], 46, 0, 0, 0);
+    set_region_offset(&region_offsets[1], 44, 0, 0, 0);
+    set_region_size(&region_sizes[0], 2, false, 0, 0, 0);
+
+    ID3D12CommandQueue_CopyTileMappings(context.queue, resource, &region_offsets[0], resource,
+            &region_offsets[1], &region_sizes[0], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    reset_command_list(context.list, context.allocator);
+
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
+    ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+    ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+    ID3D12GraphicsCommandList_SetPipelineState(context.list, check_buffer_pipeline);
+    ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+    ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+    ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < ARRAY_SIZE(buffer_region_tiles); i++)
+    {
+        set_box(&box, i, 0, 0, i + 1, 1, 1);
+        todo_if(buffer_region_tiles[i]) check_readback_data_uint(&rb.rb, &box, buffer_region_tiles[i], 0);
+    }
+
+    release_resource_readback(&rb);
+
+    ID3D12Resource_Release(resource);
+    ID3D12Heap_Release(heap);
+
+    /* Test 2D image tile mappings */
+    heap_desc.Properties = heap_properties;
+    heap_desc.Alignment = 0;
+    heap_desc.SizeInBytes = 64 * 65536;
+    heap_desc.Flags = D3D12_HEAP_FLAG_ALLOW_ONLY_NON_RT_DS_TEXTURES;
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void **)&heap);
+    ok(hr == S_OK, "Failed to create heap, hr %#x.\n", hr);
+
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
+    resource_desc.Alignment = 0;
+    resource_desc.Width = 512;
+    resource_desc.Height = 512;
+    resource_desc.DepthOrArraySize = 1;
+    resource_desc.MipLevels = 10;
+    resource_desc.Format = DXGI_FORMAT_R32_UINT;
+    resource_desc.SampleDesc.Count = 1;
+    resource_desc.SampleDesc.Quality = 0;
+    resource_desc.Layout = D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+    resource_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&resource);
+    ok(hr == S_OK, "Failed to create reserved texture, hr %#x.\n", hr);
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&resource_2);
+    ok(hr == S_OK, "Failed to create reserved texture, hr %#x.\n", hr);
+    array_resource_desc = resource_desc;
+    array_resource_desc.Width >>= 1;
+    array_resource_desc.Height >>= 1;
+    array_resource_desc.DepthOrArraySize = 3;
+    array_resource_desc.MipLevels = 1;
+    hr = ID3D12Device_CreateReservedResource(context.device, &array_resource_desc,
+            D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&array_resource);
+    ok(hr == S_OK, "Failed to create reserved texture, hr %#x.\n", hr);
+
+    num_tilings = resource_desc.MipLevels;
+    ID3D12Device_GetResourceTiling(context.device, resource, NULL, &packed_mip_info, &tile_shape, &num_tilings, 0, tilings);
+    ok(packed_mip_info.NumStandardMips >= 3, "Unexpected number of standard mips %u.\n", packed_mip_info.NumStandardMips);
+
+    srv_desc.Format = DXGI_FORMAT_R32_UINT;
+    srv_desc.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE2D;
+    srv_desc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
+    srv_desc.Texture2D.MostDetailedMip = 0;
+    srv_desc.Texture2D.MipLevels = resource_desc.MipLevels;
+    srv_desc.Texture2D.PlaneSlice = 0;
+    srv_desc.Texture2D.ResourceMinLODClamp = 0.0f;
+    ID3D12Device_CreateShaderResourceView(context.device, resource, &srv_desc, get_cpu_descriptor_handle(&context, gpu_heap, 0));
+
+    /* Map entire image */
+    tile_offsets[0] = 0;
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource,
+            1, NULL, NULL, heap, 1, NULL, tile_offsets, NULL, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    reset_command_list(context.list, context.allocator);
+
+    for (i = 0, j = 0; i < resource_desc.MipLevels; i++)
+    {
+        uav_desc.Format = DXGI_FORMAT_R32_UINT;
+        uav_desc.ViewDimension = D3D12_UAV_DIMENSION_TEXTURE2D;
+        uav_desc.Texture2D.MipSlice = i;
+        uav_desc.Texture2D.PlaneSlice = 0;
+        ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, cpu_heap, 1 + i));
+        ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, gpu_heap, 1 + i));
+
+        for (y = 0; y < max(1u, tilings[i].HeightInTiles); y++)
+        {
+            for (x = 0; x < max(1u, tilings[i].WidthInTiles); x++)
+            {
+                UINT clear_value[4] = { 0, 0, 0, 0 };
+                D3D12_RECT clear_rect;
+
+                clear_value[0] = ++j;
+                set_rect(&clear_rect, x * tile_shape.WidthInTexels, y * tile_shape.HeightInTexels,
+                        min(resource_desc.Width >> i, (x + 1) * tile_shape.WidthInTexels),
+                        min(resource_desc.Height >> i, (y + 1) * tile_shape.HeightInTexels));
+
+                ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(context.list,
+                        get_gpu_descriptor_handle(&context, gpu_heap, 1 + i),
+                        get_cpu_descriptor_handle(&context, cpu_heap, 1 + i),
+                        resource, clear_value, 1, &clear_rect);
+            }
+        }
+    }
+
+    transition_resource_state(context.list, resource, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE);
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
+    ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+    ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+    ID3D12GraphicsCommandList_SetPipelineState(context.list, check_texture_pipeline);
+    ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+    ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+    ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < j; i++)
+    {
+        set_box(&box, i, 0, 0, i + 1, 1, 1);
+        todo check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+    }
+
+    release_resource_readback(&rb);
+
+    set_region_offset(&region_offsets[0], 2, 0, 0, 0);
+    set_region_offset(&region_offsets[1], 1, 1, 0, 0);
+    set_region_offset(&region_offsets[2], 1, 1, 0, 1);
+    set_region_offset(&region_offsets[3], 0, 3, 0, 0);
+    set_region_offset(&region_offsets[4], 0, 0, 0, packed_mip_info.NumStandardMips);
+
+    set_region_size(&region_sizes[0], 3, false, 0, 0, 0);
+    set_region_size(&region_sizes[1], 4, true, 2, 2, 1);
+    set_region_size(&region_sizes[2], 2, false, 0, 0, 0);
+    set_region_size(&region_sizes[3], 4, true, 4, 1, 1);
+    set_region_size(&region_sizes[4], packed_mip_info.NumTilesForPackedMips, false, 0, 0, 0);
+
+    tile_flags[0] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[1] = D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE;
+    tile_flags[2] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[3] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[4] = D3D12_TILE_RANGE_FLAG_SKIP;
+    tile_flags[5] = D3D12_TILE_RANGE_FLAG_NONE;
+    tile_flags[6] = D3D12_TILE_RANGE_FLAG_NULL;
+    tile_flags[7] = D3D12_TILE_RANGE_FLAG_NONE;
+
+    tile_offsets[0] = 3;
+    tile_offsets[1] = 0;
+    tile_offsets[2] = 16;
+    tile_offsets[3] = 7;
+    tile_offsets[4] = 0;
+    tile_offsets[5] = 2;
+    tile_offsets[6] = 0;
+    tile_offsets[7] = 0;
+
+    tile_counts[0] = 4;
+    tile_counts[1] = 2;
+    tile_counts[2] = 3;
+    tile_counts[3] = 1;
+    tile_counts[4] = 1;
+    tile_counts[5] = 1;
+    tile_counts[6] = 1;
+    tile_counts[7] = packed_mip_info.NumTilesForPackedMips;
+
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource, 5, region_offsets, region_sizes,
+            heap, 8, tile_flags, tile_offsets, tile_counts, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    tile_offsets[7] = packed_mip_info.StartTileIndexInOverallResource;
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, resource_2, 1, &region_offsets[4], &region_sizes[4],
+            heap, 1, tile_flags, &tile_offsets[7], &tile_counts[7], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[0], 3, 1, 0, 0);
+    set_region_offset(&region_offsets[1], 1, 2, 0, 0);
+    set_region_size(&region_sizes[0], 2, true, 1, 2, 1);
+
+    ID3D12CommandQueue_CopyTileMappings(context.queue, resource, &region_offsets[0],
+            resource, &region_offsets[1], &region_sizes[0], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[0], 0, 0, 0, packed_mip_info.NumStandardMips);
+    region_offsets[1] = region_offsets[0];
+    set_region_size(&region_sizes[0], packed_mip_info.NumTilesForPackedMips, false, 0, 0, 0);
+
+    ID3D12CommandQueue_CopyTileMappings(context.queue, resource, &region_offsets[0],
+            resource_2, &region_offsets[1], &region_sizes[0], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[0], 0, 0, 0, 1);
+    set_region_offset(&region_offsets[1], 0, 0, 0, 1);
+    set_region_size(&region_sizes[0], 4, true, 2, 2, 1);
+    ID3D12CommandQueue_CopyTileMappings(context.queue, array_resource, &region_offsets[0],
+            resource, &region_offsets[1], &region_sizes[0], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_offset(&region_offsets[2], 1, 1, 0, 1);
+    set_region_offset(&region_offsets[3], 0, 0, 0, 1);
+    set_region_size(&region_sizes[1], 1, false, 0, 0, 0);
+    ID3D12CommandQueue_CopyTileMappings(context.queue, array_resource, &region_offsets[3],
+            array_resource, &region_offsets[2], &region_sizes[1], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    set_region_size(&region_sizes[0], 4, false, 0, 0, 0);
+    ID3D12CommandQueue_CopyTileMappings(context.queue, resource, &region_offsets[1],
+            array_resource, &region_offsets[0], &region_sizes[0], D3D12_TILE_MAPPING_FLAG_NONE);
+
+    reset_command_list(context.list, context.allocator);
+
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
+    ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+    ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+    ID3D12GraphicsCommandList_SetPipelineState(context.list, check_texture_pipeline);
+    ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+    ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+    ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+    transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < j; i++)
+    {
+        set_box(&box, i, 0, 0, i + 1, 1, 1);
+        todo_if(texture_region_tiles[i])
+        check_readback_data_uint(&rb.rb, &box, texture_region_tiles[i], 0);
+    }
+
+    release_resource_readback(&rb);
+
+    ID3D12Resource_Release(resource);
+    ID3D12Resource_Release(resource_2);
+    ID3D12Resource_Release(array_resource);
+
+    if (!is_amd_windows_device(context.device) && !is_radv_device(context.device))
+    {
+        /* Test 3D image tile mappings */
+        resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE3D;
+        resource_desc.Alignment = 0;
+        resource_desc.Width = 64;
+        resource_desc.Height = 64;
+        resource_desc.DepthOrArraySize = 32;
+        resource_desc.MipLevels = 2;
+        resource_desc.Format = DXGI_FORMAT_R32_UINT;
+        resource_desc.SampleDesc.Count = 1;
+        resource_desc.SampleDesc.Quality = 0;
+        resource_desc.Layout = D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+        resource_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+
+        hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+                D3D12_RESOURCE_STATE_UNORDERED_ACCESS, NULL, &IID_ID3D12Resource, (void **)&resource);
+        ok(hr == S_OK, "Failed to create reserved texture, hr %#x.\n", hr);
+
+        num_tilings = resource_desc.MipLevels;
+        ID3D12Device_GetResourceTiling(context.device, resource, NULL, &packed_mip_info, &tile_shape, &num_tilings, 0, tilings);
+        ok(packed_mip_info.NumStandardMips == 2, "Unexpected number of standard mips %u.\n", packed_mip_info.NumStandardMips);
+
+        srv_desc.Format = DXGI_FORMAT_R32_UINT;
+        srv_desc.ViewDimension = D3D12_SRV_DIMENSION_TEXTURE3D;
+        srv_desc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
+        srv_desc.Texture3D.MostDetailedMip = 0;
+        srv_desc.Texture3D.MipLevels = resource_desc.MipLevels;
+        srv_desc.Texture3D.ResourceMinLODClamp = 0.0f;
+        ID3D12Device_CreateShaderResourceView(context.device, resource, &srv_desc, get_cpu_descriptor_handle(&context, gpu_heap, 0));
+
+        /* Map entire image */
+        tile_offsets[0] = 0;
+        ID3D12CommandQueue_UpdateTileMappings(context.queue, resource,
+                1, NULL, NULL, heap, 1, NULL, tile_offsets, NULL, D3D12_TILE_MAPPING_FLAG_NONE);
+
+        reset_command_list(context.list, context.allocator);
+
+        for (i = 0, j = 0; i < resource_desc.MipLevels; i++)
+        {
+            uav_desc.Format = DXGI_FORMAT_R32_UINT;
+            uav_desc.ViewDimension = D3D12_UAV_DIMENSION_TEXTURE3D;
+            uav_desc.Texture3D.MipSlice = i;
+            uav_desc.Texture3D.FirstWSlice = 0;
+            uav_desc.Texture3D.WSize = resource_desc.DepthOrArraySize >> i;
+            ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, cpu_heap, 1 + i));
+            ID3D12Device_CreateUnorderedAccessView(context.device, resource, NULL, &uav_desc, get_cpu_descriptor_handle(&context, gpu_heap, 1 + i));
+
+            /* ClearUnorderedAccessView only takes 2D coordinates so we have to
+             * bring our own shader to initialize portions of a 3D image */
+            ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+            ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, clear_root_signature);
+            ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 1 + i));
+            ID3D12GraphicsCommandList_SetPipelineState(context.list, clear_texture_pipeline);
+
+            for (z = 0; z < max(1u, tilings[i].DepthInTiles); z++)
+            {
+                for (y = 0; y < max(1u, tilings[i].HeightInTiles); y++)
+                {
+                    for (x = 0; x < max(1u, tilings[i].WidthInTiles); x++)
+                    {
+                        UINT shader_args[4];
+                        shader_args[0] = tile_shape.WidthInTexels * x;
+                        shader_args[1] = tile_shape.HeightInTexels * y;
+                        shader_args[2] = tile_shape.DepthInTexels * z;
+                        shader_args[3] = ++j;
+
+                        ID3D12GraphicsCommandList_SetComputeRoot32BitConstants(context.list,
+                                1, ARRAY_SIZE(shader_args), shader_args, 0);
+                        ID3D12GraphicsCommandList_Dispatch(context.list, tile_shape.WidthInTexels / 4,
+                                tile_shape.HeightInTexels / 4, tile_shape.DepthInTexels / 4);
+                    }
+                }
+            }
+        }
+
+        transition_resource_state(context.list, resource, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_NON_PIXEL_SHADER_RESOURCE);
+        transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
+        ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+        ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+        ID3D12GraphicsCommandList_SetPipelineState(context.list, check_texture_3d_pipeline);
+        ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+        ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+        ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+        transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+        get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+        for (i = 0; i < j; i++)
+        {
+            set_box(&box, i, 0, 0, i + 1, 1, 1);
+            check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+        }
+
+        release_resource_readback(&rb);
+
+        set_region_offset(&region_offsets[0], 0, 0, 0, 0);
+        set_region_offset(&region_offsets[1], 0, 1, 1, 0);
+        set_region_offset(&region_offsets[2], 1, 1, 0, 0);
+        set_region_offset(&region_offsets[3], 1, 0, 0, 0);
+        set_region_offset(&region_offsets[4], 0, 1, 0, 0);
+
+        set_region_size(&region_sizes[0], 1, false, 0, 0, 0);
+        set_region_size(&region_sizes[1], 3, false, 0, 0, 0);
+        set_region_size(&region_sizes[2], 2, false, 0, 0, 0);
+        set_region_size(&region_sizes[3], 2, true,  1, 1, 2);
+        set_region_size(&region_sizes[4], 1, true,  1, 1, 1);
+
+        tile_flags[0] = D3D12_TILE_RANGE_FLAG_NONE;
+        tile_flags[1] = D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE;
+        tile_flags[2] = D3D12_TILE_RANGE_FLAG_NULL;
+
+        tile_offsets[0] = 2;
+        tile_offsets[1] = 1;
+        tile_offsets[2] = 0;
+
+        tile_counts[0] = 6;
+        tile_counts[1] = 2;
+        tile_counts[2] = 1;
+
+        ID3D12CommandQueue_UpdateTileMappings(context.queue, resource, 5, region_offsets, region_sizes,
+                heap, 3, tile_flags, tile_offsets, tile_counts, D3D12_TILE_MAPPING_FLAG_NONE);
+
+        reset_command_list(context.list, context.allocator);
+
+        transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_UNORDERED_ACCESS);
+        ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu_heap);
+        ID3D12GraphicsCommandList_SetComputeRootSignature(context.list, root_signature);
+        ID3D12GraphicsCommandList_SetPipelineState(context.list, check_texture_3d_pipeline);
+        ID3D12GraphicsCommandList_SetComputeRootDescriptorTable(context.list, 0, get_gpu_descriptor_handle(&context, gpu_heap, 0));
+        ID3D12GraphicsCommandList_SetComputeRootUnorderedAccessView(context.list, 1, readback_va);
+        ID3D12GraphicsCommandList_Dispatch(context.list, 1, 1, 1);
+        transition_resource_state(context.list, readback_buffer, D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+        get_buffer_readback_with_command_list(readback_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+        for (i = 0; i < j; i++)
+        {
+            set_box(&box, i, 0, 0, i + 1, 1, 1);
+            check_readback_data_uint(&rb.rb, &box, texture_3d_region_tiles[i], 0);
+        }
+
+        release_resource_readback(&rb);
+        ID3D12Resource_Release(resource);
+    }
+    else
+    {
+        skip("Tiles resources tier 3 not supported.\n");
+    }
+
+    ID3D12Heap_Release(heap);
+
+    ID3D12DescriptorHeap_Release(gpu_heap);
+    ID3D12DescriptorHeap_Release(cpu_heap);
+    ID3D12Resource_Release(readback_buffer);
+    ID3D12PipelineState_Release(clear_texture_pipeline);
+    ID3D12PipelineState_Release(check_texture_3d_pipeline);
+    ID3D12PipelineState_Release(check_texture_pipeline);
+    ID3D12PipelineState_Release(check_buffer_pipeline);
+    ID3D12RootSignature_Release(clear_root_signature);
+    ID3D12RootSignature_Release(root_signature);
+    destroy_test_context(&context);
+}
+
 START_TEST(d3d12)
 {
     parse_args(argc, argv);
@@ -36892,4 +37721,5 @@ START_TEST(d3d12)
     run_test(test_readback_map_stability);
     run_test(test_vs_ps_relative_addressing);
     run_test(test_get_resource_tiling);
+    run_test(test_update_tile_mappings);
 }

[PATCH 12/21] tests: Add a test for freeing underlying memory of a reserved resource.

[Conor McCarthy]

From: Conor McCarthy

As long as the reserved regions are not used, this is okay.

Based on a vkd3d-proton patch by Hans-Kristian Arntzen.
---
 tests/d3d12.c | 163 ++++++++++++++++++++++++++++++++++++++++++++++++++
 1 file changed, 163 insertions(+)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 083f9c4d5..b51bd34e4 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37541,6 +37541,168 @@ static void test_update_tile_mappings(void)
     destroy_test_context(&context);
 }

+static void test_sparse_buffer_memory_lifetime(void)
+{
+    /* Attempt to bind sparse memory, then free the underlying heap, but keep the sparse resource
+     * alive. This should confuse drivers that attempt to track BO lifetimes. */
+    D3D12_UNORDERED_ACCESS_VIEW_DESC uav_desc;
+    D3D12_SHADER_RESOURCE_VIEW_DESC srv_desc;
+    const UINT values[] = { 42, 42, 42, 42 };
+    D3D12_ROOT_PARAMETER root_parameters[2];
+    D3D12_TILE_REGION_SIZE region_size;
+    D3D12_CPU_DESCRIPTOR_HANDLE h_cpu;
+    D3D12_ROOT_SIGNATURE_DESC rs_desc;
+    D3D12_DESCRIPTOR_RANGE desc_range;
+    struct d3d12_resource_readback rb;
+    struct test_context context;
+    ID3D12DescriptorHeap *cpu;
+    ID3D12DescriptorHeap *gpu;
+    D3D12_HEAP_DESC heap_desc;
+    D3D12_RESOURCE_DESC desc;
+    ID3D12Resource *sparse;
+    ID3D12Resource *buffer;
+    ID3D12Heap *heap_live;
+    ID3D12Heap *heap;
+    unsigned int i;
+    HRESULT hr;
+
+    if (!init_compute_test_context(&context))
+        return;
+
+    memset(&rs_desc, 0, sizeof(rs_desc));
+    memset(root_parameters, 0, sizeof(root_parameters));
+    memset(&desc_range, 0, sizeof(desc_range));
+    rs_desc.NumParameters = ARRAY_SIZE(root_parameters);
+    rs_desc.pParameters = root_parameters;
+    root_parameters[0].ShaderVisibility = D3D12_SHADER_VISIBILITY_ALL;
+    root_parameters[0].ParameterType = D3D12_ROOT_PARAMETER_TYPE_UAV;
+    root_parameters[1].ShaderVisibility = D3D12_SHADER_VISIBILITY_ALL;
+    root_parameters[1].ParameterType = D3D12_ROOT_PARAMETER_TYPE_DESCRIPTOR_TABLE;
+    root_parameters[1].DescriptorTable.NumDescriptorRanges = 1;
+    root_parameters[1].DescriptorTable.pDescriptorRanges = &desc_range;
+    desc_range.NumDescriptors = 1;
+    desc_range.RangeType = D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
+    create_root_signature(context.device, &rs_desc, &context.root_signature);
+
+    memset(&heap_desc, 0, sizeof(heap_desc));
+    heap_desc.SizeInBytes = 4 * 1024 * 1024;
+    heap_desc.Properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+    heap_desc.Alignment = D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT;
+    heap_desc.Flags = D3D12_HEAP_FLAG_ALLOW_ONLY_BUFFERS;
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void**)&heap);
+    ok(SUCCEEDED(hr), "Failed to create heap, hr #%x.\n", hr);
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void**)&heap_live);
+    ok(SUCCEEDED(hr), "Failed to create heap, hr #%x.\n", hr);
+
+    memset(&desc, 0, sizeof(desc));
+    desc.Width = 64 * 1024 * 1024;
+    desc.Height = 1;
+    desc.DepthOrArraySize = 1;
+    desc.SampleDesc.Count = 1;
+    desc.Format = DXGI_FORMAT_UNKNOWN;
+    desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
+    desc.MipLevels = 1;
+    desc.Alignment = D3D12_DEFAULT_RESOURCE_PLACEMENT_ALIGNMENT;
+    desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+    desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+    hr = ID3D12Device_CreateReservedResource(context.device, &desc, D3D12_RESOURCE_STATE_UNORDERED_ACCESS,
+            NULL, &IID_ID3D12Resource, (void**)&sparse);
+    ok(SUCCEEDED(hr), "Failed to create reserved resource, hr #%x.\n", hr);
+
+    {
+        const D3D12_TILED_RESOURCE_COORDINATE region_start_coordinate = { 0 };
+        const D3D12_TILE_RANGE_FLAGS range_flag = D3D12_TILE_RANGE_FLAG_NULL;
+        const UINT offset = 0;
+        const UINT count = desc.Width / (64 * 1024);
+        region_size.UseBox = FALSE;
+        region_size.NumTiles = desc.Width / (64 * 1024);
+        ID3D12CommandQueue_UpdateTileMappings(context.queue, sparse, 1, &region_start_coordinate, &region_size,
+                NULL, 1, &range_flag, &offset, &count, D3D12_TILE_MAPPING_FLAG_NONE);
+    }
+
+    region_size.UseBox = FALSE;
+    region_size.NumTiles = 1;
+
+    for (i = 0; i < 2; i++)
+    {
+        const D3D12_TILED_RESOURCE_COORDINATE region_start_coordinate = { i, 0, 0, 0 };
+        const D3D12_TILE_RANGE_FLAGS range_flag = D3D12_TILE_RANGE_FLAG_NONE;
+        const UINT offset = i;
+        const UINT count = 1;
+
+        ID3D12CommandQueue_UpdateTileMappings(context.queue, sparse, 1, &region_start_coordinate, &region_size,
+                i ? heap_live : heap, 1, &range_flag, &offset, &count, D3D12_TILE_MAPPING_FLAG_NONE);
+    }
+    wait_queue_idle(context.device, context.queue);
+
+    buffer = create_default_buffer(context.device, 128 * 1024,
+            D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_DEST);
+    cpu = create_cpu_descriptor_heap(context.device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 1);
+    gpu = create_gpu_descriptor_heap(context.device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV, 2);
+    memset(&uav_desc, 0, sizeof(uav_desc));
+    uav_desc.ViewDimension = D3D12_UAV_DIMENSION_BUFFER;
+    uav_desc.Format = DXGI_FORMAT_R32_UINT;
+    uav_desc.Buffer.NumElements = 128 * 1024 / 4;
+    uav_desc.Buffer.FirstElement = 0;
+    ID3D12Device_CreateUnorderedAccessView(context.device, sparse, NULL, &uav_desc,
+            ID3D12DescriptorHeap_GetCPUDescriptorHandleForHeapStart(cpu));
+    ID3D12Device_CreateUnorderedAccessView(context.device, sparse, NULL, &uav_desc,
+            ID3D12DescriptorHeap_GetCPUDescriptorHandleForHeapStart(gpu));
+
+    memset(&srv_desc, 0, sizeof(srv_desc));
+    srv_desc.Buffer.FirstElement = 0;
+    srv_desc.Buffer.NumElements = 2 * 1024 * 1024;
+    srv_desc.ViewDimension = D3D12_SRV_DIMENSION_BUFFER;
+    srv_desc.Format = DXGI_FORMAT_R32_UINT;
+    srv_desc.Shader4ComponentMapping = D3D12_DEFAULT_SHADER_4_COMPONENT_MAPPING;
+
+    h_cpu = ID3D12DescriptorHeap_GetCPUDescriptorHandleForHeapStart(gpu);
+    h_cpu.ptr += ID3D12Device_GetDescriptorHandleIncrementSize(context.device, D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV);
+    ID3D12Device_CreateShaderResourceView(context.device, sparse, &srv_desc, h_cpu);
+
+    ID3D12GraphicsCommandList_SetDescriptorHeaps(context.list, 1, &gpu);
+    ID3D12GraphicsCommandList_ClearUnorderedAccessViewUint(context.list,
+            ID3D12DescriptorHeap_GetGPUDescriptorHandleForHeapStart(gpu),
+            ID3D12DescriptorHeap_GetCPUDescriptorHandleForHeapStart(cpu), sparse, values, 0, NULL);
+    transition_resource_state(context.list, sparse,
+            D3D12_RESOURCE_STATE_UNORDERED_ACCESS, D3D12_RESOURCE_STATE_COPY_SOURCE);
+    ID3D12GraphicsCommandList_CopyBufferRegion(context.list, buffer, 0, sparse, 0, 128 * 1024);
+
+    transition_resource_state(context.list, buffer, D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+    get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+    i = get_readback_uint(&rb.rb, 0, 0, 0);
+    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    i = get_readback_uint(&rb.rb, 64 * 1024 / 4, 0, 0);
+    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    release_resource_readback(&rb);
+
+    reset_command_list(context.list, context.allocator);
+
+    ID3D12Heap_Release(heap);
+
+    /* Access a resource where we can hypothetically access the freed heap memory. */
+    /* On AMD Windows native at least, if we read the freed region, we read garbage, which proves it's not required to unbind explicitly.
+     * We'd read 0 in that case. */
+    ID3D12GraphicsCommandList_CopyBufferRegion(context.list, buffer, 0, sparse, 64 * 1024, 64 * 1024);
+
+    transition_resource_state(context.list, buffer, D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    i = get_readback_uint(&rb.rb, 2048 / 4, 0, 0);
+    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    i = get_readback_uint(&rb.rb, 64 * 1024 / 4, 0, 0);
+    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    release_resource_readback(&rb);
+
+    ID3D12Resource_Release(buffer);
+    ID3D12Resource_Release(sparse);
+    ID3D12DescriptorHeap_Release(cpu);
+    ID3D12DescriptorHeap_Release(gpu);
+    ID3D12Heap_Release(heap_live);
+    destroy_test_context(&context);
+}
+
 START_TEST(d3d12)
 {
     parse_args(argc, argv);
@@ -37722,4 +37884,5 @@ START_TEST(d3d12)
     run_test(test_vs_ps_relative_addressing);
     run_test(test_get_resource_tiling);
     run_test(test_update_tile_mappings);
+    run_test(test_sparse_buffer_memory_lifetime);
 }

[PATCH 13/21] vkd3d: Pre-bind sparse texture mip tails.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/resource.c      | 101 +++++++++++++++++++++++++++++++++++++
 libs/vkd3d/vkd3d_private.h |   3 ++
 tests/d3d12.c              |   4 +-
 3 files changed, 106 insertions(+), 2 deletions(-)

diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index a6e76f197..86b8cd268 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -974,6 +974,15 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,

static void d3d12_resource_tile_info_cleanup(struct d3d12_resource *resource)
 {
+    const struct vkd3d_vk_device_procs *vk_procs = &resource->device->vk_procs;
+
+    if (!resource->tiles.subresources)
+        return;
+
+    VK_CALL(vkFreeMemory(resource->device->vk_device, resource->tiles.mip_tail_memory, NULL));
+
+    vkd3d_free(resource->tiles.bind_buffer);
+
     vkd3d_free(resource->tiles.subresources);
 }

@@ -1138,6 +1147,65 @@ void d3d12_resource_get_tiling(struct d3d12_device *device, const struct d3d12_r
     *subresource_tiling_count = i;
 }

+static void d3d12_resource_bind_sparse_mip_tail(struct d3d12_resource *resource,
+        VkSparseImageMemoryRequirements *sparse_requirements)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &resource->device->vk_procs;
+    VkSparseMemoryBind *memory_bind = resource->tiles.bind_buffer;
+    VkSparseImageOpaqueMemoryBindInfo opaque_bind_info;
+    struct d3d12_device *device = resource->device;
+    struct vkd3d_queue *vkd3d_queue;
+    VkBindSparseInfo sparse_info;
+    unsigned int i, layer_count;
+    VkDeviceSize memory_offset;
+    VkQueue vk_queue;
+    VkResult vr;
+
+    if (!resource->tiles.packed_mip_tile_count)
+        return;
+
+    vkd3d_queue = device->direct_queue;
+    if (!(vkd3d_queue->vk_queue_flags & VK_QUEUE_SPARSE_BINDING_BIT))
+    {
+        FIXME("Direct queue does not support sparse binding.\n");
+        return;
+    }
+
+    opaque_bind_info.image = resource->u.vk_image;
+    opaque_bind_info.bindCount = 1;
+    opaque_bind_info.pBinds = memory_bind;
+
+    layer_count = resource->tiles.single_mip_tail ? 1 : d3d12_resource_desc_get_layer_count(&resource->desc);
+
+    for (i = 0, memory_offset = 0; i < layer_count; ++i)
+    {
+        memory_bind->resourceOffset = sparse_requirements->imageMipTailOffset
+                + i * sparse_requirements->imageMipTailStride;
+        memory_bind->size = sparse_requirements->imageMipTailSize;
+        memory_bind->memory = resource->tiles.mip_tail_memory;
+        memory_bind->memoryOffset = memory_offset;
+        memory_bind->flags = VK_SPARSE_MEMORY_BIND_METADATA_BIT;
+        memory_offset += memory_bind->size;
+    }
+
+    memset(&sparse_info, 0, sizeof(sparse_info));
+    sparse_info.sType = VK_STRUCTURE_TYPE_BIND_SPARSE_INFO;
+    sparse_info.imageOpaqueBindCount = 1;
+    sparse_info.pImageOpaqueBinds = &opaque_bind_info;
+
+    if (!(vk_queue = vkd3d_queue_acquire(vkd3d_queue)))
+    {
+        ERR("Failed to acquire queue %p.\n", vkd3d_queue);
+        return;
+    }
+
+    if ((vr = VK_CALL(vkQueueBindSparse(vk_queue, 1, &sparse_info, VK_NULL_HANDLE))) < 0)
+        ERR("Failed to submit sparse image bind, vr %d.\n", vr);
+    /* TODO: wait on a semaphore when binding commands are implemented. */
+
+    vkd3d_queue_release(vkd3d_queue);
+}
+
 static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3d12_device *device)
 {
     unsigned int i, start_idx, subresource_count, tile_count, miplevel_idx;
@@ -1145,9 +1213,11 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
     VkSparseImageMemoryRequirements sparse_requirements_buf[3];
     VkSparseImageMemoryRequirements sparse_requirements;
     struct vkd3d_subresource_tile_info *tile_info;
+    D3D12_HEAP_PROPERTIES heap_properties;
     VkMemoryRequirements requirements;
     const VkExtent3D *tile_extent;
     uint32_t requirement_count;
+    HRESULT hr;

subresource_count = d3d12_resource_desc_get_sub_resource_count(&resource->desc);

@@ -1193,6 +1263,8 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
         }

resource->tiles.tile_extent = sparse_requirements.formatProperties.imageGranularity;
+        resource->tiles.single_mip_tail = !!(sparse_requirements.formatProperties.flags
+                & VK_SPARSE_IMAGE_FORMAT_SINGLE_MIPTAIL_BIT);
         resource->tiles.subresource_count = subresource_count;
         resource->tiles.standard_mip_count = sparse_requirements.imageMipTailSize
                 ? sparse_requirements.imageMipTailFirstLod : resource->desc.MipLevels;
@@ -1222,9 +1294,38 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
             }
         }
         resource->tiles.total_count = start_idx;
+
+        if (resource->tiles.packed_mip_tile_count)
+        {
+            memset(&heap_properties, 0, sizeof(heap_properties));
+            heap_properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+            requirements.size = sparse_requirements.imageMipTailSize;
+            if (!resource->tiles.single_mip_tail)
+                requirements.size *= d3d12_resource_desc_get_layer_count(&resource->desc);
+            if (FAILED(hr = vkd3d_allocate_device_memory(device, &heap_properties, 0, &requirements, NULL,
+                    &resource->tiles.mip_tail_memory, NULL)))
+            {
+                ERR("Failed to allocate device memory for mip tail, hr %#x.\n", hr);
+                goto error;
+            }
+        }
+
+        if (!(resource->tiles.bind_buffer = vkd3d_malloc(start_idx * max(sizeof(VkSparseImageMemoryBind),
+                sizeof(VkBufferImageCopy)))))
+        {
+            ERR("Failed to allocate binding buffer.\n");
+            goto error;
+        }
+
+        /* Vulkan requires mip tails to be always bound, while D3D12 does not, so bind them now. */
+        d3d12_resource_bind_sparse_mip_tail(resource, &sparse_requirements);
     }

return true;
+
+error:
+    d3d12_resource_tile_info_cleanup(resource);
+    return false;
 }

/* ID3D12Resource */
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 859844684..563cb4bb9 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -693,7 +693,10 @@ struct d3d12_resource_tile_info
     unsigned int standard_mip_count;
     unsigned int packed_mip_tile_count;
     unsigned int subresource_count;
+    bool single_mip_tail;
     struct vkd3d_subresource_tile_info *subresources;
+    VkDeviceMemory mip_tail_memory;
+    void *bind_buffer;
 };

#define D3D12_TILE_SIZE 0x10000u
diff --git a/tests/d3d12.c b/tests/d3d12.c
index b51bd34e4..65b6c0123 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37270,7 +37270,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+        todo_if(i < packed_mip_info.StartTileIndexInOverallResource) check_readback_data_uint(&rb.rb, &box, i + 1, 0);
     }

release_resource_readback(&rb);
@@ -37367,7 +37367,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(texture_region_tiles[i])
+        todo_if(i < packed_mip_info.StartTileIndexInOverallResource && texture_region_tiles[i])
         check_readback_data_uint(&rb.rb, &box, texture_region_tiles[i], 0);
     }

[PATCH 14/21] vkd3d: Implement ID3D12CommandQueue::UpdateTileMappings() for buffers.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 365 ++++++++++++++++++++++++++++++++++++-
 libs/vkd3d/resource.c      |  11 +-
 libs/vkd3d/vkd3d_private.h |   4 +
 tests/d3d12.c              |  13 +-
 4 files changed, 385 insertions(+), 8 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 39c0f4cd3..69451b601 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -34,7 +34,9 @@ HRESULT vkd3d_queue_create(struct d3d12_device *device,
         uint32_t family_index, const VkQueueFamilyProperties *properties, struct vkd3d_queue **queue)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkSemaphoreCreateInfo semaphore_info;
     struct vkd3d_queue *object;
+    VkResult vr;

if (!(object = vkd3d_malloc(sizeof(*object))))
         return E_OUTOFMEMORY;
@@ -54,6 +56,20 @@ HRESULT vkd3d_queue_create(struct d3d12_device *device,

memset(object->old_vk_semaphores, 0, sizeof(object->old_vk_semaphores));

+    object->tiled_binding_semaphore = VK_NULL_HANDLE;
+    if (object->vk_queue_flags & VK_QUEUE_SPARSE_BINDING_BIT)
+    {
+        semaphore_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_CREATE_INFO;
+        semaphore_info.pNext = NULL;
+        semaphore_info.flags = 0;
+        if ((vr = VK_CALL(vkCreateSemaphore(device->vk_device, &semaphore_info, NULL,
+                &object->tiled_binding_semaphore))) < 0)
+        {
+            ERR("Failed to create tiled binding semaphore, vr %d.\n", vr);
+            return hresult_from_vk_result(vr);
+        }
+    }
+
     VK_CALL(vkGetDeviceQueue(device->vk_device, family_index, 0, &object->vk_queue));

TRACE("Created queue %p for queue family index %u.\n", object, family_index);
@@ -81,6 +97,8 @@ void vkd3d_queue_destroy(struct vkd3d_queue *queue, struct d3d12_device *device)
             VK_CALL(vkDestroySemaphore(device->vk_device, queue->old_vk_semaphores[i], NULL));
     }

+    VK_CALL(vkDestroySemaphore(device->vk_device, queue->tiled_binding_semaphore, NULL));
+
     vkd3d_mutex_unlock(&queue->mutex);

vkd3d_mutex_destroy(&queue->mutex);
@@ -104,6 +122,27 @@ void vkd3d_queue_release(struct vkd3d_queue *queue)
     vkd3d_mutex_unlock(&queue->mutex);
 }

+VkResult vkd3d_queue_submit_wait_acquired(const struct vkd3d_queue *queue, VkSemaphore vk_semaphore,
+        struct d3d12_device *device)
+{
+    VkPipelineStageFlags stage_mask = VK_PIPELINE_STAGE_TOP_OF_PIPE_BIT;
+    const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+    VkSubmitInfo submit_info;
+
+    memset(&submit_info, 0, sizeof(submit_info));
+    submit_info.sType = VK_STRUCTURE_TYPE_SUBMIT_INFO;
+    submit_info.pNext = NULL;
+    submit_info.waitSemaphoreCount = 1;
+    submit_info.pWaitSemaphores = &vk_semaphore;
+    submit_info.pWaitDstStageMask = &stage_mask;
+    submit_info.commandBufferCount = 0;
+    submit_info.pCommandBuffers = 0;
+    submit_info.signalSemaphoreCount = 0;
+    submit_info.pSignalSemaphores = NULL;
+
+    return VK_CALL(vkQueueSubmit(queue->vk_queue, 1, &submit_info, VK_NULL_HANDLE));
+}
+
 static VkResult vkd3d_queue_wait_idle(struct vkd3d_queue *queue,
         const struct vkd3d_vk_device_procs *vk_procs)
 {
@@ -3887,6 +3926,104 @@ static void STDMETHODCALLTYPE d3d12_command_list_CopyResource(ID3D12GraphicsComm
     }
 }

+struct vkd3d_resource_tile_coordinate
+{
+    unsigned int x;
+    unsigned int y;
+    unsigned int z;
+};
+
+static inline unsigned int d3d12_tile_region_size_compute_tile_count(const D3D12_TILE_REGION_SIZE *region_size)
+{
+    return region_size->Width * region_size->Height * region_size->Depth;
+}
+
+static inline void d3d12_tile_region_size_set_entire_subresource(D3D12_TILE_REGION_SIZE *region_size,
+        const struct d3d12_resource *resource, unsigned int subresource)
+{
+    const struct vkd3d_tiled_region_extent *extent = &resource->tiles.subresources[subresource].extent;
+    region_size->Width = extent->width;
+    region_size->Height = extent->height;
+    region_size->Depth = extent->depth;
+}
+
+static bool resource_validate_tiled_coordinate(const struct d3d12_resource *resource,
+        const D3D12_TILED_RESOURCE_COORDINATE *coordinate)
+{
+    const struct vkd3d_tiled_region_extent *extent = &resource->tiles.subresources[coordinate->Subresource].extent;
+
+    return coordinate->Subresource < resource->tiles.subresource_count
+            && coordinate->X < extent->width && coordinate->Y < extent->height && coordinate->Z < extent->depth;
+}
+
+/* coordinate must already be validated */
+static bool resource_validate_tile_region_size(const struct d3d12_resource *resource,
+        const D3D12_TILED_RESOURCE_COORDINATE *coordinate, const D3D12_TILE_REGION_SIZE *size)
+{
+    const struct vkd3d_tiled_region_extent *extent = &resource->tiles.subresources[coordinate->Subresource].extent;
+    D3D12_TILE_REGION_SIZE max_size;
+
+    if (!size || !size->UseBox)
+        return true;
+
+    max_size.Width = extent->width - coordinate->X;
+    max_size.Height = extent->height - coordinate->Y;
+    max_size.Depth = extent->depth - coordinate->Z;
+    return size->Width <= max_size.Width && size->Height <= max_size.Height && size->Depth <= max_size.Depth;
+}
+
+/* Initialises a region in base_coordinate and region_size, where base_coordinate is always the front
+ * top left. If src_region_size->UseBox is true, start_coordinate is also the front top left, otherwise
+ * it can start anywhere within the region and the region front top left is always {0, 0, 0}. */
+static bool vkd3d_initialise_tile_region(struct vkd3d_resource_tile_coordinate *base_coordinate,
+        D3D12_TILE_REGION_SIZE *region_size, const D3D12_TILED_RESOURCE_COORDINATE *start_coordinate,
+        const D3D12_TILE_REGION_SIZE *src_region_size, const struct d3d12_resource *resource)
+{
+    unsigned int count;
+
+    if (!resource_validate_tiled_coordinate(resource, start_coordinate))
+    {
+        WARN("Invalid start coordinate (%u: %u, %u, %u).\n", start_coordinate->Subresource, start_coordinate->X,
+                start_coordinate->Y, start_coordinate->Z);
+        return false;
+    }
+    if (!resource_validate_tile_region_size(resource, start_coordinate, src_region_size))
+    {
+        WARN("Invalid region size (%u, %u, %u).\n", src_region_size->Width, src_region_size->Height,
+                src_region_size->Depth);
+        return false;
+    }
+
+    if (src_region_size)
+    {
+        *region_size = *src_region_size;
+    }
+    else
+    {
+        region_size->UseBox = false;
+        region_size->NumTiles = 1;
+    }
+
+    if (region_size->UseBox)
+    {
+        base_coordinate->x = start_coordinate->X;
+        base_coordinate->y = start_coordinate->Y;
+        base_coordinate->z = start_coordinate->Z;
+        /* NumTiles should be set by the caller. Validate it. */
+        count = d3d12_tile_region_size_compute_tile_count(region_size);
+        if (region_size->NumTiles != count)
+            WARN("NumTiles does not match the box size.\n");
+        region_size->NumTiles = count;
+    }
+    else
+    {
+        memset(base_coordinate, 0, sizeof(*base_coordinate));
+        d3d12_tile_region_size_set_entire_subresource(region_size, resource, start_coordinate->Subresource);
+    }
+
+    return true;
+}
+
 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,
@@ -6274,6 +6411,227 @@ done:
     vkd3d_mutex_unlock(&command_queue->op_mutex);
 }

+static void deaggregate_sparse_memory_bind(VkSparseBufferMemoryBindInfo *buffer_bind_info,
+        const VkSparseMemoryBind *src, unsigned int tile_count, struct d3d12_resource *resource)
+{
+    VkSparseMemoryBind *memory_binds = resource->tiles.bind_buffer;
+    unsigned int i;
+
+    for (i = 0; i < tile_count; ++i)
+    {
+        memory_binds[i].resourceOffset = src->resourceOffset + i * D3D12_TILE_SIZE;
+        memory_binds[i].size = D3D12_TILE_SIZE;
+        memory_binds[i].memory = src->memory;
+        memory_binds[i].memoryOffset = src->memoryOffset + i * D3D12_TILE_SIZE;
+        memory_binds[i].flags = src->flags;
+    }
+
+    buffer_bind_info->bindCount = tile_count;
+    buffer_bind_info->pBinds = memory_binds;
+}
+
+static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_queue *command_queue,
+        struct d3d12_resource *resource, const struct vkd3d_resource_tile_coordinate *base_coordinate,
+        D3D12_TILED_RESOURCE_COORDINATE *coordinate, const D3D12_TILE_REGION_SIZE *region_size,
+        VkDeviceMemory vk_memory, unsigned int memory_offset, unsigned int memory_tile_count, bool skip_binding)
+{
+    const struct vkd3d_vk_device_procs *vk_procs = &resource->device->vk_procs;
+    struct vkd3d_queue *vkd3d_queue = command_queue->vkd3d_queue;
+    unsigned int subresource = coordinate->Subresource;
+    VkSparseBufferMemoryBindInfo buffer_bind_info;
+    VkSparseMemoryBind memory_bind;
+    VkBindSparseInfo sparse_info;
+    unsigned int tiles_used;
+    VkResult vr;
+
+    /* The tiled resource spec for D3D11 seems to apply to D3D12 also, and states:
+     * "For mipmaps that use nonstandard tiling and/or are packed, any subresource
+     *  value that indicates any of the packed mips all refer to the same tile." */
+    if (subresource % resource->desc.MipLevels >= resource->tiles.standard_mip_count)
+    {
+        /* Already bound, but the caller expects this to use the required number of tiles,
+         * which is 1 because we bind the mip tails on resource creation and return a
+         * dummy value of 1. */
+        return 1;
+    }
+
+    memset(&sparse_info, 0, sizeof(sparse_info));
+    sparse_info.sType = VK_STRUCTURE_TYPE_BIND_SPARSE_INFO;
+
+    if (d3d12_resource_is_buffer(resource))
+    {
+        tiles_used = region_size->NumTiles;
+        tiles_used = min(tiles_used, memory_tile_count);
+
+        memory_bind.resourceOffset = coordinate->X * D3D12_TILE_SIZE;
+        coordinate->X += tiles_used;
+
+        if (skip_binding || !tiles_used)
+            return tiles_used;
+
+        memory_bind.size = tiles_used * D3D12_TILE_SIZE;
+        memory_bind.memory = vk_memory;
+        memory_bind.memoryOffset = memory_offset * D3D12_TILE_SIZE;
+        memory_bind.flags = 0;
+
+        buffer_bind_info.buffer = resource->u.vk_buffer;
+        /* A bug in NVIDIA drivers (older ones at least) requires one tile per struct to workaround. This
+         * could be skipped on other hardware by checking physical_device_info->properties2.properties.vendorID. */
+        deaggregate_sparse_memory_bind(&buffer_bind_info, &memory_bind, tiles_used, resource);
+
+        sparse_info.bufferBindCount = 1;
+        sparse_info.pBufferBinds = &buffer_bind_info;
+    }
+    else
+    {
+        vkd3d_unreachable();
+    }
+
+    sparse_info.pSignalSemaphores = &vkd3d_queue->tiled_binding_semaphore;
+    sparse_info.signalSemaphoreCount = 1;
+
+    if ((vr = VK_CALL(vkQueueBindSparse(vkd3d_queue->vk_queue, 1, &sparse_info, VK_NULL_HANDLE))) < 0)
+        ERR("Failed to submit sparse image bind, vr %d.\n", vr);
+    if (vkd3d_queue_submit_wait_acquired(vkd3d_queue, vkd3d_queue->tiled_binding_semaphore, command_queue->device) < 0)
+        ERR("Failed to submit queue wait, vr %d.\n", vr);
+
+    return tiles_used;
+}
+
+static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue *command_queue,
+        struct d3d12_resource *resource, UINT region_count,
+        const D3D12_TILED_RESOURCE_COORDINATE *region_start_coordinates,
+        const D3D12_TILE_REGION_SIZE *region_sizes,
+        struct d3d12_heap *heap,
+        UINT range_count,
+        const D3D12_TILE_RANGE_FLAGS *range_flags,
+        const UINT *heap_range_offsets,
+        const UINT *range_tile_counts,
+        D3D12_TILE_MAPPING_FLAGS flags)
+{
+    bool null_binding, aliased_binding, skip_binding, have_unsupported_aliasing;
+    VkDeviceMemory vk_memory = heap ? heap->vk_memory : VK_NULL_HANDLE;
+    unsigned int memory_offset, memory_tile_count, tiles_used;
+    struct vkd3d_resource_tile_coordinate base_coordinate;
+    D3D12_TILED_RESOURCE_COORDINATE coordinate_zero;
+    D3D12_TILE_REGION_SIZE region_size_default;
+    D3D12_TILED_RESOURCE_COORDINATE coordinate;
+    D3D12_TILE_REGION_SIZE region_size;
+    unsigned int region_idx, range_idx;
+    D3D12_TILE_RANGE_FLAGS cur_flags;
+    unsigned int tile_count_all;
+    VkQueue vk_queue;
+
+    if (d3d12_resource_is_texture(resource))
+    {
+        FIXME("Tiled textures are not implemented yet.\n");
+        return;
+    }
+
+    if (region_count == 1)
+    {
+        if (!region_sizes)
+        {
+            region_size_default.UseBox = false;
+            region_size_default.NumTiles = region_start_coordinates ? 1 : resource->tiles.total_count;
+            region_sizes = &region_size_default;
+        }
+        if (!region_start_coordinates)
+        {
+            memset(&coordinate_zero, 0, sizeof(coordinate_zero));
+            region_start_coordinates = &coordinate_zero;
+        }
+    }
+
+    if (range_count == 1 && !range_tile_counts)
+    {
+        tile_count_all = resource->tiles.total_count;
+        range_tile_counts = &tile_count_all;
+    }
+
+    if (flags)
+        WARN("Ignoring flags %#x.\n", flags);
+
+    memory_offset = heap_range_offsets ? heap_range_offsets[0] : 0;
+    memory_tile_count = range_tile_counts[0];
+    coordinate = region_start_coordinates[0];
+
+    if (!vkd3d_initialise_tile_region(&base_coordinate, &region_size, &coordinate, &region_sizes[0], resource))
+        return;
+
+    region_idx = 0;
+    range_idx = 0;
+    null_binding = false;
+    aliased_binding = false;
+    skip_binding = false;
+    have_unsupported_aliasing = false;
+
+    if (!(vk_queue = vkd3d_queue_acquire(command_queue->vkd3d_queue)))
+    {
+        ERR("Failed to acquire queue %p.\n", command_queue->vkd3d_queue);
+        return;
+    }
+
+    if (heap)
+        vkd3d_mutex_lock(&heap->mutex);
+
+    do
+    {
+        if (range_flags)
+        {
+            cur_flags = range_flags[range_idx];
+            null_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_NULL);
+            skip_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_SKIP);
+            aliased_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE);
+            if (aliased_binding && !(null_binding || skip_binding))
+            {
+                have_unsupported_aliasing = true;
+                skip_binding = true;
+            }
+            if ((cur_flags &= ~(D3D12_TILE_RANGE_FLAG_NULL | D3D12_TILE_RANGE_FLAG_SKIP | D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE)))
+                FIXME("Ignoring flags %#x.\n", cur_flags);
+        }
+
+        if (!heap_range_offsets && !null_binding)
+        {
+            WARN("Heap range offets may be NULL only if D3D12_TILE_RANGE_FLAG_NULL is used.\n");
+            break;
+        }
+
+        tiles_used = d3d12_command_queue_bind_sparse_block(command_queue, resource, &base_coordinate, &coordinate,
+                &region_size, null_binding ? VK_NULL_HANDLE : vk_memory, memory_offset,
+                aliased_binding ? 1 : memory_tile_count, skip_binding);
+
+        if (!aliased_binding)
+            memory_offset += tiles_used;
+        memory_tile_count -= tiles_used;
+        region_size.NumTiles -= tiles_used;
+
+        if (!memory_tile_count && ++range_idx < range_count)
+        {
+            memory_offset = heap_range_offsets ? heap_range_offsets[range_idx] : 0;
+            memory_tile_count = range_tile_counts[range_idx];
+        }
+
+        if (!region_size.NumTiles && ++region_idx < region_count)
+        {
+            coordinate = region_start_coordinates[region_idx];
+            if (!vkd3d_initialise_tile_region(&base_coordinate, &region_size, &coordinate,
+                    region_sizes ? &region_sizes[region_idx] : NULL, resource))
+                break;
+        }
+    }
+    while (region_idx < region_count && range_idx < range_count);
+
+    if (heap)
+        vkd3d_mutex_unlock(&heap->mutex);
+
+    if (have_unsupported_aliasing)
+        FIXME("Aliased bindings are not implemented.\n");
+
+    vkd3d_queue_release(command_queue->vkd3d_queue);
+}
+
 static void STDMETHODCALLTYPE d3d12_command_queue_CopyTileMappings(ID3D12CommandQueue *iface,
         ID3D12Resource *dst_resource,
         const D3D12_TILED_RESOURCE_COORDINATE *dst_region_start_coordinate,
@@ -7058,7 +7416,12 @@ static HRESULT d3d12_command_queue_flush_ops_locked(struct d3d12_command_queue *
                     break;

case VKD3D_CS_OP_UPDATE_MAPPINGS:
-                    FIXME("Tiled resource binding is not supported yet.\n");
+                    d3d12_command_queue_update_tile_mappings(queue, op->u.update_mappings.resource,
+                            op->u.update_mappings.region_count, op->u.update_mappings.region_start_coordinates,
+                            op->u.update_mappings.region_sizes, op->u.update_mappings.heap,
+                            op->u.update_mappings.range_count, op->u.update_mappings.range_flags,
+                            op->u.update_mappings.heap_range_offsets, op->u.update_mappings.range_tile_counts,
+                            op->u.update_mappings.flags);
                     vkd3d_free(op->u.update_mappings.region_start_coordinates);
                     vkd3d_free(op->u.update_mappings.region_sizes);
                     vkd3d_free(op->u.update_mappings.range_flags);
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index 86b8cd268..bfdc67a2d 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -1192,6 +1192,8 @@ static void d3d12_resource_bind_sparse_mip_tail(struct d3d12_resource *resource,
     sparse_info.sType = VK_STRUCTURE_TYPE_BIND_SPARSE_INFO;
     sparse_info.imageOpaqueBindCount = 1;
     sparse_info.pImageOpaqueBinds = &opaque_bind_info;
+    sparse_info.pSignalSemaphores = &vkd3d_queue->tiled_binding_semaphore;
+    sparse_info.signalSemaphoreCount = 1;

if (!(vk_queue = vkd3d_queue_acquire(vkd3d_queue)))
     {
@@ -1201,7 +1203,8 @@ static void d3d12_resource_bind_sparse_mip_tail(struct d3d12_resource *resource,

if ((vr = VK_CALL(vkQueueBindSparse(vk_queue, 1, &sparse_info, VK_NULL_HANDLE))) < 0)
         ERR("Failed to submit sparse image bind, vr %d.\n", vr);
-    /* TODO: wait on a semaphore when binding commands are implemented. */
+    if (vkd3d_queue_submit_wait_acquired(vkd3d_queue, vkd3d_queue->tiled_binding_semaphore, resource->device) < 0)
+        ERR("Failed to submit queue wait, vr %d.\n", vr);

vkd3d_queue_release(vkd3d_queue);
 }
@@ -1243,6 +1246,12 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
         resource->tiles.subresource_count = 1;
         resource->tiles.standard_mip_count = 1;
         resource->tiles.packed_mip_tile_count = 0;
+
+        if (!(resource->tiles.bind_buffer = vkd3d_malloc(resource->tiles.total_count * sizeof(VkSparseMemoryBind))))
+        {
+            ERR("Failed to allocate binding buffer.\n");
+            goto error;
+        }
     }
     else
     {
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 563cb4bb9..3eab9ed06 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1483,6 +1483,8 @@ struct vkd3d_queue
     size_t semaphore_count;

VkSemaphore old_vk_semaphores[VKD3D_MAX_VK_SYNC_OBJECTS];
+
+    VkSemaphore tiled_binding_semaphore;
 };

VkQueue vkd3d_queue_acquire(struct vkd3d_queue *queue);
@@ -1490,6 +1492,8 @@ HRESULT vkd3d_queue_create(struct d3d12_device *device, uint32_t family_index,
         const VkQueueFamilyProperties *properties, struct vkd3d_queue **queue);
 void vkd3d_queue_destroy(struct vkd3d_queue *queue, struct d3d12_device *device);
 void vkd3d_queue_release(struct vkd3d_queue *queue);
+VkResult vkd3d_queue_submit_wait_acquired(const struct vkd3d_queue *queue, VkSemaphore vk_semaphore,
+        struct d3d12_device *device);

enum vkd3d_cs_op
 {
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 65b6c0123..e670e61d1 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37080,7 +37080,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < 64; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+        check_readback_data_uint(&rb.rb, &box, i + 1, 0);
     }

release_resource_readback(&rb);
@@ -37163,7 +37163,8 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < ARRAY_SIZE(buffer_region_tiles); i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(buffer_region_tiles[i]) check_readback_data_uint(&rb.rb, &box, buffer_region_tiles[i], 0);
+        todo_if((i >= region_offsets[0].X && i < region_offsets[0].X + region_sizes[0].NumTiles) || (i >= 24 && i <= 26))
+        check_readback_data_uint(&rb.rb, &box, buffer_region_tiles[i], 0);
     }

release_resource_readback(&rb);
@@ -37671,9 +37672,9 @@ static void test_sparse_buffer_memory_lifetime(void)
     transition_resource_state(context.list, buffer, D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
     get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
     i = get_readback_uint(&rb.rb, 0, 0, 0);
-    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    ok(i == 42, "Got #%x, expected 42.\n", i);
     i = get_readback_uint(&rb.rb, 64 * 1024 / 4, 0, 0);
-    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    ok(i == 42, "Got #%x, expected 42.\n", i);
     release_resource_readback(&rb);

reset_command_list(context.list, context.allocator);
@@ -37690,9 +37691,9 @@ static void test_sparse_buffer_memory_lifetime(void)
     get_buffer_readback_with_command_list(buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);

i = get_readback_uint(&rb.rb, 2048 / 4, 0, 0);
-    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    ok(i == 42, "Got #%x, expected 42.\n", i);
     i = get_readback_uint(&rb.rb, 64 * 1024 / 4, 0, 0);
-    todo ok(i == 42, "Got #%x, expected 42.\n", i);
+    ok(i == 42, "Got #%x, expected 42.\n", i);
     release_resource_readback(&rb);

ID3D12Resource_Release(buffer);

[PATCH 15/21] vkd3d: Implement ID3D12CommandQueue::UpdateTileMappings() for textures.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 162 +++++++++++++++++++++++++++++++++++--
 libs/vkd3d/vkd3d_private.h |  13 +++
 tests/d3d12.c              |   4 +-
 3 files changed, 169 insertions(+), 10 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 69451b601..a70d1b01f 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4024,6 +4024,41 @@ static bool vkd3d_initialise_tile_region(struct vkd3d_resource_tile_coordinate *
     return true;
 }

+static bool d3d12_tiled_resource_coordinate_normalise(const struct vkd3d_resource_tile_coordinate *base_coordinate,
+        const D3D12_TILE_REGION_SIZE *region_extent, D3D12_TILED_RESOURCE_COORDINATE *coordinate)
+{
+    unsigned int carry;
+
+    /* This should compile branchless on most hardware. */
+    carry = coordinate->X >= base_coordinate->x + region_extent->Width;
+    coordinate->Y += carry;
+    coordinate->X -= region_extent->Width & -carry;
+
+    carry = coordinate->Y >= base_coordinate->y + region_extent->Height;
+    coordinate->Z += carry;
+    coordinate->Y -= region_extent->Height & -carry;
+
+    carry = coordinate->Z >= base_coordinate->z + region_extent->Depth;
+    coordinate->Subresource += carry;
+    coordinate->Z -= region_extent->Depth & -carry;
+
+    return carry;
+}
+
+static inline void vk_offset_convert_tiles_to_texels(VkOffset3D *offset, const VkExtent3D *tile_extent)
+{
+    offset->x *= tile_extent->width;
+    offset->y *= tile_extent->height;
+    offset->z *= tile_extent->depth;
+}
+
+static inline void vk_extent_convert_tiles_to_texels(VkExtent3D *extent, const VkExtent3D *tile_extent)
+{
+    extent->width *= tile_extent->width;
+    extent->height *= tile_extent->height;
+    extent->depth *= tile_extent->depth;
+}
+
 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,
@@ -6411,6 +6446,61 @@ done:
     vkd3d_mutex_unlock(&command_queue->op_mutex);
 }

+static unsigned int vkd3d_set_sparse_image_bind_region(VkSparseImageMemoryBind *memory_bind,
+        const struct vkd3d_resource_tile_coordinate *base_coordinate,
+        D3D12_TILED_RESOURCE_COORDINATE *coordinate,
+        const D3D12_TILE_REGION_SIZE *region_extent, const struct d3d12_resource *resource,
+        unsigned int memory_tile_count)
+{
+    unsigned int height, depth, remaining, max_tile_count, tile_count, layer_stride;
+    VkOffset3D *offset = &memory_bind->offset;
+    VkExtent3D *extent = &memory_bind->extent;
+    bool partial_x, partial_y;
+
+    max_tile_count = min(region_extent->NumTiles, memory_tile_count);
+    tile_count = max_tile_count;
+    partial_x = coordinate->X > base_coordinate->x;
+    partial_y = coordinate->Y > base_coordinate->y;
+
+    offset->x = coordinate->X;
+    offset->y = coordinate->Y;
+    offset->z = coordinate->Z;
+
+    /* Grab the largest possible width */
+    remaining = region_extent->Width - (coordinate->X - base_coordinate->x);
+    extent->width = min(remaining, tile_count);
+    coordinate->X += extent->width;
+    tile_count -= extent->width;
+    extent->height = 1;
+    extent->depth = 1;
+    if (d3d12_tiled_resource_coordinate_normalise(base_coordinate, region_extent, coordinate)
+            || partial_x || !tile_count || !(height = tile_count / region_extent->Width))
+        goto done;
+
+    /* Expand the height */
+    remaining = region_extent->Height - (coordinate->Y - base_coordinate->y);
+    remaining = min(remaining, height);
+    extent->height += remaining;
+    coordinate->Y += remaining;
+    tile_count -= region_extent->Width * remaining;
+    if (d3d12_tiled_resource_coordinate_normalise(base_coordinate, region_extent, coordinate)
+            || partial_y || !tile_count
+            || !(depth = tile_count / (layer_stride = region_extent->Width * region_extent->Height)))
+        goto done;
+
+    /* Expand the depth */
+    remaining = region_extent->Depth - (coordinate->Z - base_coordinate->z);
+    remaining = min(remaining, depth);
+    extent->depth += remaining;
+    coordinate->Z += remaining;
+    tile_count -= layer_stride * remaining;
+
+    d3d12_tiled_resource_coordinate_normalise(base_coordinate, region_extent, coordinate);
+
+done:
+    return max_tile_count - tile_count;
+}
+
 static void deaggregate_sparse_memory_bind(VkSparseBufferMemoryBindInfo *buffer_bind_info,
         const VkSparseMemoryBind *src, unsigned int tile_count, struct d3d12_resource *resource)
 {
@@ -6430,6 +6520,40 @@ static void deaggregate_sparse_memory_bind(VkSparseBufferMemoryBindInfo *buffer_
     buffer_bind_info->pBinds = memory_binds;
 }

+static void deaggregate_sparse_image_memory_bind(VkSparseImageMemoryBindInfo *image_bind_info,
+        const VkSparseImageMemoryBind *src, struct d3d12_resource *resource)
+{
+    VkSparseImageMemoryBind *image_memory_binds = resource->tiles.bind_buffer;
+    const VkExtent3D *tile_extent = &resource->tiles.tile_extent;
+    unsigned int i, x, y, z, tile_count;
+
+    tile_count = src->extent.width * src->extent.height * src->extent.depth;
+
+    for (z = 0, i = 0; z < src->extent.depth; ++z)
+    {
+        for (y = 0; y < src->extent.height; ++y)
+        {
+            for (x = 0; x < src->extent.width; ++x, ++i)
+            {
+                image_memory_binds[i].subresource = src->subresource;
+                image_memory_binds[i].offset.x = src->offset.x + x;
+                image_memory_binds[i].offset.y = src->offset.y + y;
+                image_memory_binds[i].offset.z = src->offset.z + z;
+                vk_offset_convert_tiles_to_texels(&image_memory_binds[i].offset, tile_extent);
+                image_memory_binds[i].extent.width = tile_extent->width;
+                image_memory_binds[i].extent.height = tile_extent->height;
+                image_memory_binds[i].extent.depth = tile_extent->depth;
+                image_memory_binds[i].memory = src->memory;
+                image_memory_binds[i].memoryOffset = src->memoryOffset + i * D3D12_TILE_SIZE;
+                image_memory_binds[i].flags = src->flags;
+            }
+        }
+    }
+
+    image_bind_info->bindCount = tile_count;
+    image_bind_info->pBinds = image_memory_binds;
+}
+
 static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_queue *command_queue,
         struct d3d12_resource *resource, const struct vkd3d_resource_tile_coordinate *base_coordinate,
         D3D12_TILED_RESOURCE_COORDINATE *coordinate, const D3D12_TILE_REGION_SIZE *region_size,
@@ -6439,6 +6563,8 @@ static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_q
     struct vkd3d_queue *vkd3d_queue = command_queue->vkd3d_queue;
     unsigned int subresource = coordinate->Subresource;
     VkSparseBufferMemoryBindInfo buffer_bind_info;
+    VkSparseImageMemoryBindInfo image_bind_info;
+    VkSparseImageMemoryBind image_memory_bind;
     VkSparseMemoryBind memory_bind;
     VkBindSparseInfo sparse_info;
     unsigned int tiles_used;
@@ -6484,7 +6610,25 @@ static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_q
     }
     else
     {
-        vkd3d_unreachable();
+        d3d12_resource_get_vk_subresource(resource, subresource, &image_memory_bind.subresource);
+
+        tiles_used = vkd3d_set_sparse_image_bind_region(&image_memory_bind,
+                base_coordinate, coordinate, region_size, resource, memory_tile_count);
+
+        if (skip_binding || !tiles_used)
+            return tiles_used;
+
+        image_bind_info.image = resource->u.vk_image;
+
+        image_memory_bind.memory = vk_memory;
+        image_memory_bind.memoryOffset = memory_offset * D3D12_TILE_SIZE;
+        image_memory_bind.flags = 0;
+
+        /* NVIDIA bug (see above).*/
+        deaggregate_sparse_image_memory_bind(&image_bind_info, &image_memory_bind, resource);
+
+        sparse_info.imageBindCount = 1;
+        sparse_info.pImageBinds = &image_bind_info;
     }

sparse_info.pSignalSemaphores = &vkd3d_queue->tiled_binding_semaphore;
@@ -6510,8 +6654,8 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
         D3D12_TILE_MAPPING_FLAGS flags)
 {
     bool null_binding, aliased_binding, skip_binding, have_unsupported_aliasing;
+    unsigned int memory_offset, memory_tile_count, tiles_used, subresource;
     VkDeviceMemory vk_memory = heap ? heap->vk_memory : VK_NULL_HANDLE;
-    unsigned int memory_offset, memory_tile_count, tiles_used;
     struct vkd3d_resource_tile_coordinate base_coordinate;
     D3D12_TILED_RESOURCE_COORDINATE coordinate_zero;
     D3D12_TILE_REGION_SIZE region_size_default;
@@ -6522,12 +6666,6 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
     unsigned int tile_count_all;
     VkQueue vk_queue;

-    if (d3d12_resource_is_texture(resource))
-    {
-        FIXME("Tiled textures are not implemented yet.\n");
-        return;
-    }
-
     if (region_count == 1)
     {
         if (!region_sizes)
@@ -6559,6 +6697,7 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
     if (!vkd3d_initialise_tile_region(&base_coordinate, &region_size, &coordinate, &region_sizes[0], resource))
         return;

+    subresource = coordinate.Subresource;
     region_idx = 0;
     range_idx = 0;
     null_binding = false;
@@ -6577,6 +6716,13 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue

do
     {
+        if (coordinate.Subresource != subresource)
+        {
+            if ((subresource = coordinate.Subresource) >= resource->tiles.subresource_count)
+                break;
+            d3d12_tile_region_size_set_entire_subresource(&region_size, resource, subresource);
+        }
+
         if (range_flags)
         {
             cur_flags = range_flags[range_idx];
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 3eab9ed06..dfc49e866 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1865,6 +1865,19 @@ 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 void d3d12_resource_get_vk_subresource(const struct d3d12_resource *resource, unsigned int subresource,
+        VkImageSubresource *vk_subresource)
+{
+    const struct vkd3d_format *format = resource->format;
+    const D3D12_RESOURCE_DESC *desc = &resource->desc;
+
+    assert(format->plane_count == 1);
+
+    vk_subresource->mipLevel = subresource % desc->MipLevels;
+    vk_subresource->arrayLayer = subresource / desc->MipLevels;
+    vk_subresource->aspectMask = format->vk_aspect_mask;
+}
+
 static inline unsigned int vkd3d_compute_workgroup_count(unsigned int thread_count, unsigned int workgroup_size)
 {
     return (thread_count + workgroup_size - 1) / workgroup_size;
diff --git a/tests/d3d12.c b/tests/d3d12.c
index e670e61d1..b30615b4b 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37271,7 +37271,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(i < packed_mip_info.StartTileIndexInOverallResource) check_readback_data_uint(&rb.rb, &box, i + 1, 0);
+        check_readback_data_uint(&rb.rb, &box, i + 1, 0);
     }

release_resource_readback(&rb);
@@ -37368,7 +37368,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(i < packed_mip_info.StartTileIndexInOverallResource && texture_region_tiles[i])
+        todo_if(i == 6 || i == 7 || i == 9 || i == 11 || i == 16)
         check_readback_data_uint(&rb.rb, &box, texture_region_tiles[i], 0);
     }

[PATCH 16/21] vkd3d: Support aliased tile binding if available.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 5 +++--
 libs/vkd3d/device.c        | 1 +
 libs/vkd3d/resource.c      | 4 ++++
 libs/vkd3d/vkd3d_private.h | 1 +
 tests/d3d12.c              | 4 ++--
 5 files changed, 11 insertions(+), 4 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index a70d1b01f..38934a5d6 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -6654,6 +6654,7 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
         D3D12_TILE_MAPPING_FLAGS flags)
 {
     bool null_binding, aliased_binding, skip_binding, have_unsupported_aliasing;
+    bool can_alias = command_queue->device->vk_info.sparse_residency_aliased;
     unsigned int memory_offset, memory_tile_count, tiles_used, subresource;
     VkDeviceMemory vk_memory = heap ? heap->vk_memory : VK_NULL_HANDLE;
     struct vkd3d_resource_tile_coordinate base_coordinate;
@@ -6729,7 +6730,7 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
             null_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_NULL);
             skip_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_SKIP);
             aliased_binding = !!(cur_flags & D3D12_TILE_RANGE_FLAG_REUSE_SINGLE_TILE);
-            if (aliased_binding && !(null_binding || skip_binding))
+            if (aliased_binding && !(null_binding || skip_binding) && !can_alias)
             {
                 have_unsupported_aliasing = true;
                 skip_binding = true;
@@ -6773,7 +6774,7 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
         vkd3d_mutex_unlock(&heap->mutex);

if (have_unsupported_aliasing)
-        FIXME("Aliased bindings are not implemented.\n");
+        FIXME("Aliased bindings are not supported by the device.\n");

vkd3d_queue_release(command_queue->vkd3d_queue);
 }
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index b1eba634f..3c54632ca 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -1465,6 +1465,7 @@ static HRESULT vkd3d_init_device_caps(struct d3d12_device *device,
     vulkan_info->device_limits = physical_device_info->properties2.properties.limits;
     vulkan_info->sparse_properties = physical_device_info->properties2.properties.sparseProperties;
     vulkan_info->sparse_residency_3d = features->sparseResidencyImage3D;
+    vulkan_info->sparse_residency_aliased = features->sparseResidencyAliased;
     vulkan_info->rasterization_stream = physical_device_info->xfb_properties.transformFeedbackRasterizationStreamSelect;
     vulkan_info->transform_feedback_queries = physical_device_info->xfb_properties.transformFeedbackQueries;
     vulkan_info->uav_read_without_format = features->shaderStorageImageReadWithoutFormat;
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index bfdc67a2d..fc5c30242 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -667,6 +667,8 @@ HRESULT vkd3d_create_buffer(struct d3d12_device *device,
         buffer_info.flags |= VK_BUFFER_CREATE_SPARSE_BINDING_BIT;
         if (device->vk_info.sparse_properties.residencyNonResidentStrict)
             buffer_info.flags |= VK_BUFFER_CREATE_SPARSE_RESIDENCY_BIT;
+        if (device->vk_info.sparse_residency_aliased)
+            buffer_info.flags |= VK_BUFFER_CREATE_SPARSE_ALIASED_BIT;
     }

buffer_info.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT
@@ -828,6 +830,8 @@ static HRESULT vkd3d_create_image(struct d3d12_device *device,
         image_info.flags |= VK_IMAGE_CREATE_SPARSE_BINDING_BIT;
         if (device->vk_info.sparse_properties.residencyNonResidentStrict)
             image_info.flags |= VK_IMAGE_CREATE_SPARSE_RESIDENCY_BIT;
+        if (device->vk_info.sparse_residency_aliased)
+            image_info.flags |= VK_IMAGE_CREATE_SPARSE_ALIASED_BIT;
     }

image_info.imageType = vk_image_type_from_d3d12_resource_dimension(desc->Dimension);
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index dfc49e866..53aa1d102 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -151,6 +151,7 @@ struct vkd3d_vulkan_info

VkPhysicalDeviceSparseProperties sparse_properties;
     bool sparse_residency_3d;
+    bool sparse_residency_aliased;

VkPhysicalDeviceTexelBufferAlignmentPropertiesEXT texel_buffer_alignment_properties;

diff --git a/tests/d3d12.c b/tests/d3d12.c
index b30615b4b..0b9cadf74 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37163,7 +37163,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < ARRAY_SIZE(buffer_region_tiles); i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if((i >= region_offsets[0].X && i < region_offsets[0].X + region_sizes[0].NumTiles) || (i >= 24 && i <= 26))
+        todo_if(i >= region_offsets[0].X && i < region_offsets[0].X + region_sizes[0].NumTiles)
         check_readback_data_uint(&rb.rb, &box, buffer_region_tiles[i], 0);
     }

@@ -37368,7 +37368,7 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(i == 6 || i == 7 || i == 9 || i == 11 || i == 16)
+        todo_if(i == 7 || i == 11 || i == 16)
         check_readback_data_uint(&rb.rb, &box, texture_region_tiles[i], 0);
     }

[PATCH 17/21] vkd3d: Implement ID3D12CommandQueue::CopyTileMappings() for buffers.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c       | 131 ++++++++++++++++++++++++++++++++++++-
 libs/vkd3d/resource.c      |  17 +++++
 libs/vkd3d/vkd3d_private.h |   8 +++
 tests/d3d12.c              |   1 -
 4 files changed, 155 insertions(+), 2 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index 38934a5d6..ec306cd04 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -6501,6 +6501,27 @@ done:
     return max_tile_count - tile_count;
 }

+static void d3d12_resource_update_buffer_tile_mappings(struct d3d12_resource *resource,
+        const VkSparseMemoryBind *memory_bind)
+{
+    struct vkd3d_subresource_tile_mapping *mappings = resource->tiles.subresources[0].mappings;
+    VkDeviceSize offset = memory_bind->memoryOffset;
+    unsigned int i, end;
+
+    vkd3d_mutex_lock(&resource->tiles.mutex);
+
+    i = memory_bind->resourceOffset / D3D12_TILE_SIZE;
+    end = i + memory_bind->size / D3D12_TILE_SIZE;
+    for (; i < end; ++i)
+    {
+        mappings[i].vk_memory = memory_bind->memory;
+        mappings[i].byte_offset = offset;
+        offset += D3D12_TILE_SIZE;
+    }
+
+    vkd3d_mutex_unlock(&resource->tiles.mutex);
+}
+
 static void deaggregate_sparse_memory_bind(VkSparseBufferMemoryBindInfo *buffer_bind_info,
         const VkSparseMemoryBind *src, unsigned int tile_count, struct d3d12_resource *resource)
 {
@@ -6600,6 +6621,8 @@ static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_q
         memory_bind.memoryOffset = memory_offset * D3D12_TILE_SIZE;
         memory_bind.flags = 0;

+        d3d12_resource_update_buffer_tile_mappings(resource, &memory_bind);
+
         buffer_bind_info.buffer = resource->u.vk_buffer;
         /* A bug in NVIDIA drivers (older ones at least) requires one tile per struct to workaround. This
          * could be skipped on other hardware by checking physical_device_info->properties2.properties.vendorID. */
@@ -6779,6 +6802,109 @@ static void d3d12_command_queue_update_tile_mappings(struct d3d12_command_queue
     vkd3d_queue_release(command_queue->vkd3d_queue);
 }

+static void vkd3d_copy_vk_tile_mapping_region(const struct d3d12_command_queue *command_queue,
+        const struct vkd3d_resource_tile_coordinate *dst_base, D3D12_TILED_RESOURCE_COORDINATE *dst_loc,
+        const D3D12_TILE_REGION_SIZE *dst_extent, struct d3d12_resource *dst_resource,
+        const struct vkd3d_resource_tile_coordinate *src_base, D3D12_TILED_RESOURCE_COORDINATE *src_loc,
+        const D3D12_TILE_REGION_SIZE *src_extent, const struct d3d12_resource *src_resource)
+{
+    const struct vkd3d_subresource_tile_info *src_info = &src_resource->tiles.subresources[src_loc->Subresource];
+    const struct vkd3d_vk_device_procs *vk_procs = &dst_resource->device->vk_procs;
+    struct vkd3d_queue *vkd3d_queue = command_queue->vkd3d_queue;
+    const struct vkd3d_subresource_tile_mapping *src_mapping;
+    VkSparseBufferMemoryBindInfo buffer_bind_info;
+    VkSparseMemoryBind *memory_bind;
+    VkBindSparseInfo sparse_info;
+    unsigned int i, src_idx;
+    VkResult vr;
+
+    src_mapping = src_info->mappings;
+
+    memset(&sparse_info, 0, sizeof(sparse_info));
+    sparse_info.sType = VK_STRUCTURE_TYPE_BIND_SPARSE_INFO;
+
+    if (d3d12_resource_is_buffer(dst_resource))
+    {
+        sparse_info.bufferBindCount = 1;
+        sparse_info.pBufferBinds = &buffer_bind_info;
+        memory_bind = dst_resource->tiles.bind_buffer;
+        buffer_bind_info.buffer = dst_resource->u.vk_buffer;
+        buffer_bind_info.bindCount = src_extent->NumTiles;
+        buffer_bind_info.pBinds = memory_bind;
+
+        src_idx = src_loc->X;
+        for (i = 0; i < src_extent->NumTiles; ++i, ++src_idx)
+        {
+            memory_bind[i].resourceOffset = dst_loc->X * dst_resource->tiles.tile_extent.width;
+            memory_bind[i].size = dst_resource->tiles.tile_extent.width;
+            memory_bind[i].memory = src_mapping[src_idx].vk_memory;
+            memory_bind[i].memoryOffset = src_mapping[src_idx].byte_offset;
+            memory_bind[i].flags = 0;
+            ++dst_loc->X;
+        }
+
+        for (i = 0; i < src_extent->NumTiles; ++i)
+            d3d12_resource_update_buffer_tile_mappings(dst_resource, &memory_bind[i]);
+    }
+    else
+    {
+        vkd3d_unreachable();
+    }
+
+    sparse_info.pSignalSemaphores = &vkd3d_queue->tiled_binding_semaphore;
+    sparse_info.signalSemaphoreCount = 1;
+    if ((vr = VK_CALL(vkQueueBindSparse(vkd3d_queue->vk_queue, 1, &sparse_info, VK_NULL_HANDLE))) < 0)
+        ERR("Failed to submit sparse image bind, vr %d.\n", vr);
+    if (vkd3d_queue_submit_wait_acquired(vkd3d_queue, vkd3d_queue->tiled_binding_semaphore, command_queue->device) < 0)
+        ERR("Failed to submit queue wait, vr %d.\n", vr);
+}
+
+static void d3d12_command_queue_copy_tile_mappings(struct d3d12_command_queue *command_queue,
+        struct d3d12_resource *dst_resource,
+        const D3D12_TILED_RESOURCE_COORDINATE *dst_region_start_coordinate,
+        struct d3d12_resource *src_resource,
+        const D3D12_TILED_RESOURCE_COORDINATE *src_region_start_coordinate,
+        const D3D12_TILE_REGION_SIZE *region_size,
+        D3D12_TILE_MAPPING_FLAGS flags)
+{
+    struct vkd3d_resource_tile_coordinate dst_base, src_base;
+    D3D12_TILED_RESOURCE_COORDINATE dst_loc, src_loc;
+    D3D12_TILE_REGION_SIZE dst_extent, src_extent;
+    VkQueue vk_queue;
+
+    if (d3d12_resource_is_texture(dst_resource) || d3d12_resource_is_texture(src_resource))
+    {
+        FIXME("Not implemented for textures.\n");
+        return;
+    }
+
+    dst_loc = *dst_region_start_coordinate;
+    src_loc = *src_region_start_coordinate;
+    if (!vkd3d_initialise_tile_region(&dst_base, &dst_extent, &dst_loc, region_size, dst_resource)
+            || !vkd3d_initialise_tile_region(&src_base, &src_extent, &src_loc, region_size, src_resource))
+    {
+        WARN("Invalid tile region.\n");
+        return;
+    }
+
+    if (!src_extent.NumTiles || (dst_resource == src_resource
+            && dst_loc.X == src_loc.X
+            && dst_loc.Y == src_loc.Y
+            && dst_loc.Z == src_loc.Z && dst_loc.Subresource == src_loc.Subresource))
+        return;
+
+    if (!(vk_queue = vkd3d_queue_acquire(command_queue->vkd3d_queue)))
+    {
+        ERR("Failed to acquire queue %p.\n", command_queue->vkd3d_queue);
+        return;
+    }
+
+    vkd3d_copy_vk_tile_mapping_region(command_queue, &dst_base, &dst_loc, &dst_extent, dst_resource,
+            &src_base, &src_loc, &src_extent, src_resource);
+
+    vkd3d_queue_release(command_queue->vkd3d_queue);
+}
+
 static void STDMETHODCALLTYPE d3d12_command_queue_CopyTileMappings(ID3D12CommandQueue *iface,
         ID3D12Resource *dst_resource,
         const D3D12_TILED_RESOURCE_COORDINATE *dst_region_start_coordinate,
@@ -7577,7 +7703,10 @@ static HRESULT d3d12_command_queue_flush_ops_locked(struct d3d12_command_queue *
                     break;

case VKD3D_CS_OP_COPY_MAPPINGS:
-                    FIXME("Tiled resource mapping copying is not supported yet.\n");
+                    d3d12_command_queue_copy_tile_mappings(queue, op->u.copy_mappings.dst_resource,
+                            &op->u.copy_mappings.dst_region_start_coordinate, op->u.copy_mappings.src_resource,
+                            &op->u.copy_mappings.src_region_start_coordinate, &op->u.copy_mappings.region_size,
+                            op->u.copy_mappings.flags);
                     break;

default:
diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c
index fc5c30242..8d3d3a7fb 100644
--- a/libs/vkd3d/resource.c
+++ b/libs/vkd3d/resource.c
@@ -979,14 +979,19 @@ HRESULT vkd3d_get_image_allocation_info(struct d3d12_device *device,
 static void d3d12_resource_tile_info_cleanup(struct d3d12_resource *resource)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &resource->device->vk_procs;
+    unsigned int i;

if (!resource->tiles.subresources)
         return;

+    vkd3d_mutex_destroy(&resource->tiles.mutex);
+
     VK_CALL(vkFreeMemory(resource->device->vk_device, resource->tiles.mip_tail_memory, NULL));

vkd3d_free(resource->tiles.bind_buffer);

+    for (i = 0; i < resource->tiles.subresource_count; ++i)
+        vkd3d_free(resource->tiles.subresources[i].mappings);
     vkd3d_free(resource->tiles.subresources);
 }

@@ -1234,6 +1239,8 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
         return false;
     }

+    vkd3d_mutex_init(&resource->tiles.mutex);
+
     if (d3d12_resource_is_buffer(resource))
     {
         tile_info = &resource->tiles.subresources[0];
@@ -1256,6 +1263,11 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
             ERR("Failed to allocate binding buffer.\n");
             goto error;
         }
+        if (!(tile_info->mappings = vkd3d_calloc(resource->tiles.total_count, sizeof(*tile_info->mappings))))
+        {
+            ERR("Failed to allocate mapping buffer.\n");
+            goto error;
+        }
     }
     else
     {
@@ -1298,6 +1310,11 @@ static bool d3d12_resource_init_tiles(struct d3d12_resource *resource, struct d3
             {
                 tile_count = tile_info->extent.width * tile_info->extent.height * tile_info->extent.depth;
                 start_idx += tile_count;
+                if (!(tile_info->mappings = vkd3d_calloc(tile_count, sizeof(*tile_info->mappings))))
+                {
+                    ERR("Failed to allocate mapping buffer.\n");
+                    goto error;
+                }
                 tile_info->count = tile_count;
             }
             else if (miplevel_idx == resource->tiles.standard_mip_count)
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 53aa1d102..c999ef79e 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -680,11 +680,18 @@ struct vkd3d_tiled_region_extent
     unsigned int depth;
 };

+struct vkd3d_subresource_tile_mapping
+{
+    VkDeviceMemory vk_memory;
+    VkDeviceSize byte_offset;
+};
+
 struct vkd3d_subresource_tile_info
 {
     unsigned int offset;
     unsigned int count;
     struct vkd3d_tiled_region_extent extent;
+    struct vkd3d_subresource_tile_mapping *mappings;
 };

struct d3d12_resource_tile_info
@@ -695,6 +702,7 @@ struct d3d12_resource_tile_info
     unsigned int packed_mip_tile_count;
     unsigned int subresource_count;
     bool single_mip_tail;
+    struct vkd3d_mutex mutex;
     struct vkd3d_subresource_tile_info *subresources;
     VkDeviceMemory mip_tail_memory;
     void *bind_buffer;
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 0b9cadf74..b1e68fa1f 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37163,7 +37163,6 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < ARRAY_SIZE(buffer_region_tiles); i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(i >= region_offsets[0].X && i < region_offsets[0].X + region_sizes[0].NumTiles)
         check_readback_data_uint(&rb.rb, &box, buffer_region_tiles[i], 0);
     }

[PATCH 18/21] vkd3d: Implement ID3D12CommandQueue::CopyTileMappings() for textures.

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c | 107 +++++++++++++++++++++++++++++++++++++++----
 tests/d3d12.c        |   1 -
 2 files changed, 97 insertions(+), 11 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index ec306cd04..d61ee7177 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -6522,6 +6522,39 @@ static void d3d12_resource_update_buffer_tile_mappings(struct d3d12_resource *re
     vkd3d_mutex_unlock(&resource->tiles.mutex);
 }

+/* Never called for mip tails. */
+static void d3d12_resource_update_image_tile_mappings(struct d3d12_resource *resource, unsigned int subresource,
+        const VkSparseImageMemoryBind *image_memory_bind)
+{
+    struct vkd3d_subresource_tile_mapping *layer, *column, *row;
+    VkDeviceSize offset = image_memory_bind->memoryOffset;
+    const struct vkd3d_tiled_region_extent *extent;
+    unsigned int x, y, z, layer_stride;
+
+    extent = &resource->tiles.subresources[subresource].extent;
+
+    vkd3d_mutex_lock(&resource->tiles.mutex);
+
+    layer_stride = extent->width * extent->height;
+    layer = &resource->tiles.subresources[subresource].mappings[image_memory_bind->offset.z * layer_stride];
+    for (z = 0; z < image_memory_bind->extent.depth; ++z, layer += layer_stride)
+    {
+        column = layer + image_memory_bind->offset.y * extent->width;
+        for (y = 0; y < image_memory_bind->extent.height; ++y, column += extent->width)
+        {
+            row = column + image_memory_bind->offset.x;
+            for (x = 0; x < image_memory_bind->extent.width; ++x)
+            {
+                row[x].vk_memory = image_memory_bind->memory;
+                row[x].byte_offset = offset;
+                offset += D3D12_TILE_SIZE;
+            }
+        }
+    }
+
+    vkd3d_mutex_unlock(&resource->tiles.mutex);
+}
+
 static void deaggregate_sparse_memory_bind(VkSparseBufferMemoryBindInfo *buffer_bind_info,
         const VkSparseMemoryBind *src, unsigned int tile_count, struct d3d12_resource *resource)
 {
@@ -6647,6 +6680,8 @@ static unsigned int d3d12_command_queue_bind_sparse_block(struct d3d12_command_q
         image_memory_bind.memoryOffset = memory_offset * D3D12_TILE_SIZE;
         image_memory_bind.flags = 0;

+        d3d12_resource_update_image_tile_mappings(resource, subresource, &image_memory_bind);
+
         /* NVIDIA bug (see above).*/
         deaggregate_sparse_image_memory_bind(&image_bind_info, &image_memory_bind, resource);

@@ -6813,12 +6848,16 @@ static void vkd3d_copy_vk_tile_mapping_region(const struct d3d12_command_queue *
     struct vkd3d_queue *vkd3d_queue = command_queue->vkd3d_queue;
     const struct vkd3d_subresource_tile_mapping *src_mapping;
     VkSparseBufferMemoryBindInfo buffer_bind_info;
+    VkSparseImageMemoryBindInfo image_bind_info;
+    VkSparseImageMemoryBind *image_memory_bind;
+    unsigned int i, src_idx, subresource;
     VkSparseMemoryBind *memory_bind;
     VkBindSparseInfo sparse_info;
-    unsigned int i, src_idx;
+    bool is_end;
     VkResult vr;

src_mapping = src_info->mappings;
+    subresource = dst_loc->Subresource;

memset(&sparse_info, 0, sizeof(sparse_info));
     sparse_info.sType = VK_STRUCTURE_TYPE_BIND_SPARSE_INFO;
@@ -6848,7 +6887,42 @@ static void vkd3d_copy_vk_tile_mapping_region(const struct d3d12_command_queue *
     }
     else
     {
-        vkd3d_unreachable();
+        sparse_info.imageBindCount = 1;
+        sparse_info.pImageBinds = &image_bind_info;
+        image_memory_bind = dst_resource->tiles.bind_buffer;
+        image_bind_info.image = dst_resource->u.vk_image;
+        image_bind_info.bindCount = src_extent->NumTiles;
+        image_bind_info.pBinds = image_memory_bind;
+
+        for (i = 0, is_end = false; i < src_extent->NumTiles && !is_end; ++i)
+        {
+            src_idx = src_loc->X + src_loc->Y * src_info->extent.width
+                    + src_loc->Z * src_info->extent.width * src_info->extent.height;
+            d3d12_resource_get_vk_subresource(dst_resource, dst_loc->Subresource, &image_memory_bind[i].subresource);
+            image_memory_bind[i].offset.x = dst_loc->X;
+            image_memory_bind[i].offset.y = dst_loc->Y;
+            image_memory_bind[i].offset.z = dst_loc->Z;
+            image_memory_bind[i].extent.width = 1;
+            image_memory_bind[i].extent.height = 1;
+            image_memory_bind[i].extent.depth = 1;
+            image_memory_bind[i].memory = src_mapping[src_idx].vk_memory;
+            image_memory_bind[i].memoryOffset = src_mapping[src_idx].byte_offset;
+            image_memory_bind[i].flags = 0;
+            ++src_loc->X;
+            is_end = d3d12_tiled_resource_coordinate_normalise(src_base, src_extent, src_loc);
+            ++dst_loc->X;
+            is_end |= d3d12_tiled_resource_coordinate_normalise(dst_base, dst_extent, dst_loc);
+        }
+        if (i < src_extent->NumTiles && src_loc->Subresource < src_resource->tiles.subresource_count
+                && dst_loc->Subresource < dst_resource->tiles.subresource_count)
+            FIXME("Multiple sub-resource support is not implemented.\n");
+
+        for (i = 0; i < src_extent->NumTiles; ++i)
+        {
+            d3d12_resource_update_image_tile_mappings(dst_resource, subresource, &image_memory_bind[i]);
+            vk_offset_convert_tiles_to_texels(&image_memory_bind[i].offset, &dst_resource->tiles.tile_extent);
+            vk_extent_convert_tiles_to_texels(&image_memory_bind[i].extent, &dst_resource->tiles.tile_extent);
+        }
     }

sparse_info.pSignalSemaphores = &vkd3d_queue->tiled_binding_semaphore;
@@ -6870,14 +6944,9 @@ static void d3d12_command_queue_copy_tile_mappings(struct d3d12_command_queue *c
     struct vkd3d_resource_tile_coordinate dst_base, src_base;
     D3D12_TILED_RESOURCE_COORDINATE dst_loc, src_loc;
     D3D12_TILE_REGION_SIZE dst_extent, src_extent;
+    bool dst_miptail, src_miptail;
     VkQueue vk_queue;

-    if (d3d12_resource_is_texture(dst_resource) || d3d12_resource_is_texture(src_resource))
-    {
-        FIXME("Not implemented for textures.\n");
-        return;
-    }
-
     dst_loc = *dst_region_start_coordinate;
     src_loc = *src_region_start_coordinate;
     if (!vkd3d_initialise_tile_region(&dst_base, &dst_extent, &dst_loc, region_size, dst_resource)
@@ -6893,14 +6962,32 @@ static void d3d12_command_queue_copy_tile_mappings(struct d3d12_command_queue *c
             && dst_loc.Z == src_loc.Z && dst_loc.Subresource == src_loc.Subresource))
         return;

+    dst_miptail = dst_loc.Subresource % dst_resource->desc.MipLevels >= dst_resource->tiles.standard_mip_count;
+    src_miptail = src_loc.Subresource % src_resource->desc.MipLevels >= src_resource->tiles.standard_mip_count;
+    if (dst_miptail != src_miptail)
+    {
+        /* This scenario makes little sense. */
+        FIXME("Not implemented for copies between a mip tail and a standard mip.\n");
+        return;
+    }
+
     if (!(vk_queue = vkd3d_queue_acquire(command_queue->vkd3d_queue)))
     {
         ERR("Failed to acquire queue %p.\n", command_queue->vkd3d_queue);
         return;
     }

-    vkd3d_copy_vk_tile_mapping_region(command_queue, &dst_base, &dst_loc, &dst_extent, dst_resource,
-            &src_base, &src_loc, &src_extent, src_resource);
+    if (dst_miptail)
+    {
+        /* Already bound. */
+        vkd3d_queue_release(command_queue->vkd3d_queue);
+        return;
+    }
+    else
+    {
+        vkd3d_copy_vk_tile_mapping_region(command_queue, &dst_base, &dst_loc, &dst_extent, dst_resource,
+                &src_base, &src_loc, &src_extent, src_resource);
+    }

vkd3d_queue_release(command_queue->vkd3d_queue);
 }
diff --git a/tests/d3d12.c b/tests/d3d12.c
index b1e68fa1f..3f9f712ce 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -37367,7 +37367,6 @@ static void test_update_tile_mappings(void)
     for (i = 0; i < j; i++)
     {
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo_if(i == 7 || i == 11 || i == 16)
         check_readback_data_uint(&rb.rb, &box, texture_region_tiles[i], 0);
     }

[PATCH 19/21] tests/d3d12: Add tests for CopyTiles().

[Conor McCarthy]

From: Conor McCarthy

Based on a vkd3d-proton patch by Philip Rebohle.
---
 tests/d3d12.c | 249 +++++++++++++++++++++++++++++++++++++++++++++++++-
 1 file changed, 247 insertions(+), 2 deletions(-)

diff --git a/tests/d3d12.c b/tests/d3d12.c
index 3f9f712ce..17089673f 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -192,8 +192,7 @@ static void get_buffer_readback_with_command_list(ID3D12Resource *buffer, DXGI_F
     resource_desc.Flags = D3D12_RESOURCE_FLAG_DENY_SHADER_RESOURCE;

hr = ID3D12Resource_GetHeapProperties(buffer, &heap_properties, NULL);
-    ok(SUCCEEDED(hr), "Failed to get heap properties.\n");
-    if (heap_properties.Type == D3D12_HEAP_TYPE_READBACK)
+    if (SUCCEEDED(hr) && heap_properties.Type == D3D12_HEAP_TYPE_READBACK)
     {
         rb_buffer = buffer;
         ID3D12Resource_AddRef(rb_buffer);
@@ -37702,6 +37701,251 @@ static void test_sparse_buffer_memory_lifetime(void)
     destroy_test_context(&context);
 }

+static void test_copy_tiles(void)
+{
+#define TILE_SIZE 65536
+    ID3D12Resource *tiled_resource, *dst_buffer, *src_buffer;
+    D3D12_TILED_RESOURCE_COORDINATE region_offset;
+    uint32_t tile_offset, buffer_offset;
+    D3D12_TILE_REGION_SIZE region_size;
+    D3D12_RESOURCE_DESC resource_desc;
+    struct d3d12_resource_readback rb;
+    struct test_context_desc desc;
+    struct test_context context;
+    D3D12_HEAP_DESC heap_desc;
+    uint32_t *buffer_data;
+    unsigned int i, x, y;
+    ID3D12Heap *heap;
+    D3D12_BOX box;
+    HRESULT hr;
+
+    static const struct
+    {
+        uint32_t x;
+        uint32_t y;
+        uint32_t tile_idx;
+    }
+    image_tiles[] =
+    {
+        {1, 0, 0}, {2, 0, 1}, {1, 1, 2}, {2, 1, 3},
+        {3, 1, 4}, {0, 2, 5}, {1, 2, 6},
+    };
+
+    memset(&desc, 0, sizeof(desc));
+    desc.rt_width = 640;
+    desc.rt_height = 480;
+    desc.rt_format = DXGI_FORMAT_R8G8B8A8_UNORM;
+    if (!init_test_context(&context, &desc))
+        return;
+
+    memset(&heap_desc, 0, sizeof(heap_desc));
+    heap_desc.Properties.Type = D3D12_HEAP_TYPE_DEFAULT;
+    heap_desc.SizeInBytes = TILE_SIZE * 16;
+
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
+    resource_desc.Alignment = 0;
+    resource_desc.Width = heap_desc.SizeInBytes;
+    resource_desc.Height = 1;
+    resource_desc.DepthOrArraySize = 1;
+    resource_desc.MipLevels = 1;
+    resource_desc.Format = DXGI_FORMAT_UNKNOWN;
+    resource_desc.SampleDesc.Count = 1;
+    resource_desc.SampleDesc.Quality = 0;
+    resource_desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+    resource_desc.Flags = D3D12_RESOURCE_FLAG_NONE;
+    hr = ID3D12Device_CreateCommittedResource(context.device, &heap_desc.Properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COPY_DEST, NULL, &IID_ID3D12Resource, (void **)&src_buffer);
+    ok(hr == S_OK, "Failed to create buffer, hr %#x.\n", hr);
+    hr = ID3D12Device_CreateCommittedResource(context.device, &heap_desc.Properties, D3D12_HEAP_FLAG_NONE,
+            &resource_desc, D3D12_RESOURCE_STATE_COPY_DEST, NULL, &IID_ID3D12Resource, (void **)&dst_buffer);
+    ok(hr == S_OK, "Failed to create buffer, hr %#x.\n", hr);
+
+    buffer_data = malloc(resource_desc.Width);
+    for (i = 0; i < resource_desc.Width / sizeof(*buffer_data); i++)
+        buffer_data[i] = i;
+    upload_buffer_data(src_buffer, 0, resource_desc.Width, buffer_data, context.queue, context.list);
+
+    reset_command_list(context.list, context.allocator);
+    transition_resource_state(context.list, src_buffer,
+            D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    /* Test buffer */
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_COPY_DEST, NULL, &IID_ID3D12Resource, (void **)&tiled_resource);
+    ok(hr == S_OK, "Failed to create tiled buffer, hr %#x.\n", hr);
+
+    heap_desc.Flags = D3D12_HEAP_FLAG_ALLOW_ONLY_BUFFERS;
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void **)&heap);
+    ok(hr == S_OK, "Failed to create heap, hr %#x.\n", hr);
+
+    tile_offset = 0;
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, tiled_resource,
+            1, NULL, NULL, heap, 1, NULL, &tile_offset, NULL, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    /* Copy source tiles 0-2 with a 32-byte offset to buffer tiles 4-6 */
+    set_region_offset(&region_offset, 4, 0, 0, 0);
+    set_region_size(&region_size, 3, false, 0, 0, 0);
+
+    buffer_offset = 32;
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            src_buffer, buffer_offset, D3D12_TILE_COPY_FLAG_LINEAR_BUFFER_TO_SWIZZLED_TILED_RESOURCE);
+
+    transition_resource_state(context.list, tiled_resource,
+            D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(tiled_resource, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < 3 * TILE_SIZE / sizeof(*buffer_data); i += 1024)
+    {
+        uint32_t offset = i + 4 * TILE_SIZE / sizeof(*buffer_data);
+        set_box(&box, offset, 0, 0, offset + 1, 1, 1);
+        todo check_readback_data_uint(&rb.rb, &box, buffer_data[i + buffer_offset / sizeof(*buffer_data)], 0);
+    }
+
+    release_resource_readback(&rb);
+
+    reset_command_list(context.list, context.allocator);
+
+    /* Read tiles 5-6 from the tiled resource */
+    set_region_offset(&region_offset, 5, 0, 0, 0);
+    set_region_size(&region_size, 1, false, 0, 0, 0);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            dst_buffer, 0, D3D12_TILE_COPY_FLAG_SWIZZLED_TILED_RESOURCE_TO_LINEAR_BUFFER);
+
+    /* NONE behaves the same as SWIZZLED_TILED_RESOURCE_TO_LINEAR_BUFFER */
+    set_region_offset(&region_offset, 6, 0, 0, 0);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            dst_buffer, TILE_SIZE, D3D12_TILE_COPY_FLAG_NONE);
+
+    transition_resource_state(context.list, dst_buffer,
+            D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(dst_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < 2 * TILE_SIZE / sizeof(*buffer_data); i += 1024)
+    {
+        uint32_t offset = i + (TILE_SIZE + buffer_offset) / sizeof(*buffer_data);
+        set_box(&box, i, 0, 0, i + 1, 1, 1);
+        todo check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+    }
+
+    release_resource_readback(&rb);
+
+    ID3D12Resource_Release(tiled_resource);
+    ID3D12Heap_Release(heap);
+
+    /* Test image */
+    resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
+    resource_desc.Alignment = 0;
+    resource_desc.Width = 512;
+    resource_desc.Height = 512;
+    resource_desc.DepthOrArraySize = 1;
+    resource_desc.MipLevels = 1;
+    resource_desc.Format = DXGI_FORMAT_R32_UINT;
+    resource_desc.SampleDesc.Count = 1;
+    resource_desc.SampleDesc.Quality = 0;
+    resource_desc.Layout = D3D12_TEXTURE_LAYOUT_64KB_UNDEFINED_SWIZZLE;
+    resource_desc.Flags = D3D12_RESOURCE_FLAG_NONE;
+
+    hr = ID3D12Device_CreateReservedResource(context.device, &resource_desc,
+            D3D12_RESOURCE_STATE_COPY_DEST, NULL, &IID_ID3D12Resource, (void **)&tiled_resource);
+    ok(hr == S_OK, "Failed to create tiled buffer, hr %#x.\n", hr);
+
+    heap_desc.Flags = D3D12_HEAP_FLAG_ALLOW_ONLY_NON_RT_DS_TEXTURES;
+    hr = ID3D12Device_CreateHeap(context.device, &heap_desc, &IID_ID3D12Heap, (void **)&heap);
+    ok(hr == S_OK, "Failed to create heap, hr %#x.\n", hr);
+
+    tile_offset = 0;
+    ID3D12CommandQueue_UpdateTileMappings(context.queue, tiled_resource,
+            1, NULL, NULL, heap, 1, NULL, &tile_offset, NULL, D3D12_TILE_MAPPING_FLAG_NONE);
+
+    reset_command_list(context.list, context.allocator);
+
+    /* Copy source tiles 0-3 to 2x2 region at (1,0) */
+    set_region_offset(&region_offset, 1, 0, 0, 0);
+    set_region_size(&region_size, 4, true, 2, 2, 1);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            src_buffer, 0, D3D12_TILE_COPY_FLAG_LINEAR_BUFFER_TO_SWIZZLED_TILED_RESOURCE);
+
+    /* Copy source tiles 4-6 to (3,1), (0,2) and (1,2) */
+    set_region_offset(&region_offset, 3, 1, 0, 0);
+    set_region_size(&region_size, 3, false, 0, 0, 0);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            src_buffer, 4 * TILE_SIZE, D3D12_TILE_COPY_FLAG_LINEAR_BUFFER_TO_SWIZZLED_TILED_RESOURCE);
+
+    transition_resource_state(context.list, tiled_resource,
+            D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_resource_readback_with_command_list(tiled_resource, 0, &rb, context.queue, context.list);
+
+    for (i = 0; i < ARRAY_SIZE(image_tiles); i++)
+    {
+        for (y = 0; y < 128; y += 32)
+        {
+            for (x = 0; x < 128; x += 32)
+            {
+                uint32_t offset = image_tiles[i].tile_idx * TILE_SIZE / sizeof(*buffer_data) + 128 * y + x;
+                set_box(&box, 128 * image_tiles[i].x + x, 128 * image_tiles[i].y + y, 0,
+                        128 * image_tiles[i].x + x + 1, 128 * image_tiles[i].y + y + 1, 1);
+                todo_if(buffer_data[offset]) check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+            }
+        }
+    }
+
+    release_resource_readback(&rb);
+
+    reset_command_list(context.list, context.allocator);
+
+    transition_resource_state(context.list, dst_buffer,
+            D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_COPY_DEST);
+
+    /* Read 0-3 to 2x2 region at (1,0) */
+    set_region_offset(&region_offset, 1, 0, 0, 0);
+    set_region_size(&region_size, 4, true, 2, 2, 1);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            dst_buffer, 0, D3D12_TILE_COPY_FLAG_SWIZZLED_TILED_RESOURCE_TO_LINEAR_BUFFER);
+
+    /* Read tiles (3,1), (0,2) and (1,2) */
+    set_region_offset(&region_offset, 3, 1, 0, 0);
+    set_region_size(&region_size, 3, false, 0, 0, 0);
+
+    ID3D12GraphicsCommandList_CopyTiles(context.list, tiled_resource, &region_offset, &region_size,
+            dst_buffer, 4 * TILE_SIZE, D3D12_TILE_COPY_FLAG_NONE);
+
+    transition_resource_state(context.list, dst_buffer,
+            D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE);
+
+    get_buffer_readback_with_command_list(dst_buffer, DXGI_FORMAT_R32_UINT, &rb, context.queue, context.list);
+
+    for (i = 0; i < ARRAY_SIZE(image_tiles); i++)
+    {
+        for (x = 0; x < TILE_SIZE / sizeof(uint32_t); x += 1024)
+        {
+            uint32_t offset = image_tiles[i].tile_idx * TILE_SIZE / sizeof(uint32_t) + x;
+            set_box(&box, offset, 0, 0, offset + 1, 1, 1);
+            todo_if(buffer_data[offset]) check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+        }
+    }
+
+    release_resource_readback(&rb);
+
+    ID3D12Resource_Release(tiled_resource);
+    ID3D12Heap_Release(heap);
+
+    ID3D12Resource_Release(src_buffer);
+    ID3D12Resource_Release(dst_buffer);
+
+    free(buffer_data);
+    destroy_test_context(&context);
+#undef TILE_SIZE
+}
+
 START_TEST(d3d12)
 {
     parse_args(argc, argv);
@@ -37884,4 +38128,5 @@ START_TEST(d3d12)
     run_test(test_get_resource_tiling);
     run_test(test_update_tile_mappings);
     run_test(test_sparse_buffer_memory_lifetime);
+    run_test(test_copy_tiles);
 }

[PATCH 20/21] vkd3d: Introduce a config flag to report tiled resources as unsupported.

[Conor McCarthy]

From: Conor McCarthy

---
 README                     | 4 ++++
 libs/vkd3d/device.c        | 4 ++++
 libs/vkd3d/vkd3d_private.h | 1 +
 3 files changed, 9 insertions(+)

diff --git a/README b/README
index 465d5f915..87784033e 100644
--- a/README
+++ b/README
@@ -51,6 +51,10 @@ commas or semicolons.
    even when the output supports colour.

* VKD3D_CONFIG - a list of options that change the behavior of libvkd3d.
+    * tiled_tier_0 - Do not report support for tiled resources. This option may
+      prevent failures due to unsupported sparse image formats which can occur
+      with some hardware/drivers. Tiled resources are still supported if an app
+      uses them anyway and the Vulkan driver supports them.
     * virtual_heaps - Create descriptors for each D3D12 root signature
       descriptor range instead of entire descriptor heaps. Useful when push
       constant or bound descriptor limits are exceeded.
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 3c54632ca..fce248919 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -498,6 +498,7 @@ static const struct vkd3d_debug_option vkd3d_config_options[] =
 {
     {"virtual_heaps", VKD3D_CONFIG_FLAG_VIRTUAL_HEAPS}, /* always use virtual descriptor heaps */
     {"vk_debug", VKD3D_CONFIG_FLAG_VULKAN_DEBUG}, /* enable Vulkan debug extensions */
+    {"tiled_tier_0", VKD3D_CONFIG_FLAG_TILED_TIER_0}, /* do not report tiled resource support */
 };

static uint64_t vkd3d_init_config_flags(void)
@@ -1485,6 +1486,9 @@ static HRESULT vkd3d_init_device_caps(struct d3d12_device *device,
     else
         device->feature_options.TiledResourcesTier = D3D12_TILED_RESOURCES_TIER_3;

+    if (device->vkd3d_instance->config_flags & VKD3D_CONFIG_FLAG_TILED_TIER_0)
+        device->feature_options.TiledResourcesTier = D3D12_TILED_RESOURCES_TIER_NOT_SUPPORTED;
+
     /* FIXME: Implement tiled resources. */
     if (device->feature_options.TiledResourcesTier)
     {
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index c999ef79e..0964be49c 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -165,6 +165,7 @@ enum vkd3d_config_flags
 {
     VKD3D_CONFIG_FLAG_VULKAN_DEBUG = 0x00000001,
     VKD3D_CONFIG_FLAG_VIRTUAL_HEAPS = 0x00000002,
+    VKD3D_CONFIG_FLAG_TILED_TIER_0 = 0x00000004,
 };

struct vkd3d_instance

[PATCH 21/21] vkd3d: Implement ID3D12GraphicsCommandList::CopyTiles().

[Conor McCarthy]

From: Conor McCarthy

---
 libs/vkd3d/command.c | 136 ++++++++++++++++++++++++++++++++++++++++++-
 libs/vkd3d/device.c  |   7 ---
 tests/d3d12.c        |  46 +++++++++++----
 3 files changed, 169 insertions(+), 20 deletions(-)

diff --git a/libs/vkd3d/command.c b/libs/vkd3d/command.c
index d61ee7177..f149c64d1 100644
--- a/libs/vkd3d/command.c
+++ b/libs/vkd3d/command.c
@@ -4059,15 +4059,145 @@ static inline void vk_extent_convert_tiles_to_texels(VkExtent3D *extent, const V
     extent->depth *= tile_extent->depth;
 }

+static void d3d12_command_list_copy_texture_tiles(struct d3d12_command_list *list,
+        struct d3d12_resource *resource, const struct d3d12_resource *buffer, uint64_t buffer_offset,
+        const D3D12_TILED_RESOURCE_COORDINATE *tile_region_start_coordinate,
+        const D3D12_TILE_REGION_SIZE *tile_region_size, bool to_image)
+{
+    D3D12_TILED_RESOURCE_COORDINATE coordinate = *tile_region_start_coordinate;
+    VkBufferImageCopy *vk_copies = resource->tiles.bind_buffer;
+    struct vkd3d_subresource_tile_mapping *mapping = NULL;
+    struct vkd3d_resource_tile_coordinate base_coordinate;
+    const struct vkd3d_tiled_region_extent *extent;
+    const struct vkd3d_vk_device_procs *vk_procs;
+    VkImageSubresourceLayers vk_subresource;
+    D3D12_TILE_REGION_SIZE region_size;
+    unsigned int i, j, subresource;
+    VkImageLayout layout;
+
+    vk_procs = &list->device->vk_procs;
+    layout = to_image ? VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL : VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL;
+
+    subresource = 0;
+
+    vkd3d_mutex_lock(&resource->tiles.mutex);
+
+    /* In RADV at least, only one tile can be copied per VkBufferImageCopy, maybe because of
+     * a bug in address calculations for multiple buffer tiles. */
+    for (i = 0, j = 0; i < tile_region_size->NumTiles; ++i, ++mapping)
+    {
+        if (subresource != coordinate.Subresource || !i)
+        {
+            if ((subresource = coordinate.Subresource) >= resource->tiles.subresource_count)
+                break;
+            if (subresource % resource->desc.MipLevels >= resource->tiles.standard_mip_count)
+            {
+                /* The spec doesn't state we should skip over packed mips and copy the next layer, so just break. */
+                WARN("CopyTiles() does not support packed mips.\n");
+                break;
+            }
+
+            if (!vkd3d_initialise_tile_region(&base_coordinate, &region_size, &coordinate, tile_region_size, resource))
+            {
+                WARN("Invalid tile region.\n");
+                break;
+            }
+
+            vk_image_subresource_layers_from_d3d12(&vk_subresource, resource->format,
+                    subresource, resource->desc.MipLevels);
+            extent = &resource->tiles.subresources[subresource].extent;
+            mapping = &resource->tiles.subresources[subresource].mappings[coordinate.X
+                    + coordinate.Y * extent->width + coordinate.Z * extent->width * extent->height];
+        }
+
+        if (mapping->vk_memory)
+        {
+            vk_copies[j].bufferOffset = buffer_offset;
+            vk_copies[j].bufferRowLength = resource->tiles.tile_extent.width;
+            vk_copies[j].bufferImageHeight = resource->tiles.tile_extent.height;
+
+            vk_copies[j].imageSubresource = vk_subresource;
+            vk_copies[j].imageOffset.x = coordinate.X;
+            vk_copies[j].imageOffset.y = coordinate.Y;
+            vk_copies[j].imageOffset.z = coordinate.Z;
+            vk_offset_convert_tiles_to_texels(&vk_copies[j].imageOffset, &resource->tiles.tile_extent);
+            vk_copies[j++].imageExtent = resource->tiles.tile_extent;
+        }
+
+        buffer_offset += D3D12_TILE_SIZE;
+        ++coordinate.X;
+        d3d12_tiled_resource_coordinate_normalise(&base_coordinate, &region_size, &coordinate);
+    }
+
+    if (to_image)
+        VK_CALL(vkCmdCopyBufferToImage(list->vk_command_buffer, buffer->u.vk_buffer, resource->u.vk_image, layout, j, vk_copies));
+    else
+        VK_CALL(vkCmdCopyImageToBuffer(list->vk_command_buffer, resource->u.vk_image, layout, buffer->u.vk_buffer, j, vk_copies));
+
+    vkd3d_mutex_unlock(&resource->tiles.mutex);
+}
+
 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,
+        const D3D12_TILE_REGION_SIZE *tile_region_size, ID3D12Resource *buffer_iface, UINT64 buffer_offset,
         D3D12_TILE_COPY_FLAGS flags)
 {
-    FIXME("iface %p, tiled_resource %p, tile_region_start_coordinate %p, tile_region_size %p, "
-            "buffer %p, buffer_offset %#"PRIx64", flags %#x stub!\n",
+    struct d3d12_resource *resource = unsafe_impl_from_ID3D12Resource(tiled_resource);
+    struct d3d12_resource *buffer = unsafe_impl_from_ID3D12Resource(buffer_iface);
+    struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList2(iface);
+    const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
+    D3D12_TILE_COPY_FLAGS ignored_flags;
+    bool to_resource;
+
+    TRACE("iface %p, tiled_resource %p, tile_region_start_coordinate %p, tile_region_size %p, "
+            "buffer %p, buffer_offset %#"PRIx64", flags %#x.\n",
             iface, tiled_resource, tile_region_start_coordinate, tile_region_size,
             buffer, buffer_offset, flags);
+
+    if (!resource || !buffer)
+        return;
+
+    d3d12_command_list_track_resource_usage(list, resource);
+
+    d3d12_command_list_end_current_render_pass(list);
+
+    if (tile_region_start_coordinate->Subresource >= resource->tiles.subresource_count)
+    {
+        WARN("Invalid sub resource %u.\n", tile_region_start_coordinate->Subresource);
+        return;
+    }
+
+    ignored_flags = flags;
+    flags &= D3D12_TILE_COPY_FLAG_LINEAR_BUFFER_TO_SWIZZLED_TILED_RESOURCE |
+            D3D12_TILE_COPY_FLAG_SWIZZLED_TILED_RESOURCE_TO_LINEAR_BUFFER;
+    if (vkd3d_popcount(flags) > 1)
+    {
+        WARN("Invalid flags %#x. Skipping.\n", flags);
+        return;
+    }
+    if ((ignored_flags ^= flags))
+        WARN("Ignoring flags %#x.\n", ignored_flags);
+    /* No flags defaults to resource-to-buffer. */
+    to_resource = flags == D3D12_TILE_COPY_FLAG_LINEAR_BUFFER_TO_SWIZZLED_TILED_RESOURCE;
+
+    if (d3d12_resource_is_texture(resource))
+    {
+        d3d12_command_list_copy_texture_tiles(list, resource, buffer, buffer_offset,
+                tile_region_start_coordinate, tile_region_size, to_resource);
+    }
+    else
+    {
+        VkDeviceSize tiled_offset = tile_region_start_coordinate->X * D3D12_TILE_SIZE;
+        VkBufferCopy buffer_copy;
+
+        buffer_copy.srcOffset = to_resource ? buffer_offset : tiled_offset;
+        buffer_copy.dstOffset = to_resource ? tiled_offset : buffer_offset;
+        buffer_copy.size = tile_region_size->NumTiles * D3D12_TILE_SIZE;
+
+        VK_CALL(vkCmdCopyBuffer(list->vk_command_buffer,
+                to_resource ? buffer->u.vk_buffer : resource->u.vk_buffer,
+                to_resource ? resource->u.vk_buffer : buffer->u.vk_buffer, 1, &buffer_copy));
+    }
 }

static void STDMETHODCALLTYPE d3d12_command_list_ResolveSubresource(ID3D12GraphicsCommandList2 *iface,
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index fce248919..c8da7cb22 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -1489,13 +1489,6 @@ static HRESULT vkd3d_init_device_caps(struct d3d12_device *device,
     if (device->vkd3d_instance->config_flags & VKD3D_CONFIG_FLAG_TILED_TIER_0)
         device->feature_options.TiledResourcesTier = D3D12_TILED_RESOURCES_TIER_NOT_SUPPORTED;

-    /* FIXME: Implement tiled resources. */
-    if (device->feature_options.TiledResourcesTier)
-    {
-        WARN("Tiled resources are not implemented yet.\n");
-        device->feature_options.TiledResourcesTier = D3D12_TILED_RESOURCES_TIER_NOT_SUPPORTED;
-    }
-
     if (device->vk_info.device_limits.maxPerStageDescriptorSamplers <= 16)
         device->feature_options.ResourceBindingTier = D3D12_RESOURCE_BINDING_TIER_1;
     else if (device->vk_info.device_limits.maxPerStageDescriptorUniformBuffers <= 14)
diff --git a/tests/d3d12.c b/tests/d3d12.c
index 17089673f..2bc7d91e8 100644
--- a/tests/d3d12.c
+++ b/tests/d3d12.c
@@ -36446,6 +36446,7 @@ static uint32_t compute_tile_count(uint32_t resource_size, uint32_t mip, uint32_
 static void test_get_resource_tiling(void)
 {
     D3D12_SUBRESOURCE_TILING tilings_alt[17];
+    D3D12_TILED_RESOURCES_TIER tiled_tier;
     D3D12_PACKED_MIP_INFO packed_mip_info;
     D3D12_SUBRESOURCE_TILING tilings[17];
     UINT num_resource_tiles, num_tilings;
@@ -36455,7 +36456,6 @@ static void test_get_resource_tiling(void)
     D3D12_TILE_SHAPE tile_shape;
     ID3D12Resource *resource;
     unsigned int i, j;
-    bool no_tier_3;
     HRESULT hr;

static const struct
@@ -36555,6 +36555,13 @@ static void test_get_resource_tiling(void)
     if (!init_test_context(&context, &desc))
         return;

+    if ((tiled_tier = get_tiled_resources_tier(context.device)) < D3D12_TILED_RESOURCES_TIER_1)
+    {
+        skip("Tiled resources not supported by device.\n");
+        destroy_test_context(&context);
+        return;
+    }
+
     /* Test behaviour with various parameter combinations */
     resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D;
     resource_desc.Alignment = 0;
@@ -36603,16 +36610,13 @@ static void test_get_resource_tiling(void)

ID3D12Resource_Release(resource);

-    /* Tiled tier is not included in feature support yet so as not to break future bisections. */
-    no_tier_3 = is_amd_windows_device(context.device) || is_radv_device(context.device);
-
     /* Test actual tiling properties */
     for (i = 0; i < ARRAY_SIZE(tests); i++, vkd3d_test_pop_context())
     {
         unsigned int tile_index = 0;
         vkd3d_test_push_context("test %u", i);

-        if (no_tier_3 && tests[i].min_tier > D3D12_TILED_RESOURCES_TIER_2)
+        if (tests[i].min_tier > tiled_tier)
         {
             skip("Tiled resources tier %u not supported.\n", tests[i].min_tier);
             continue;
@@ -36748,6 +36752,7 @@ static void test_update_tile_mappings(void)
     D3D12_TILE_REGION_SIZE region_sizes[8];
     D3D12_GPU_VIRTUAL_ADDRESS readback_va;
     D3D12_HEAP_PROPERTIES heap_properties;
+    D3D12_TILED_RESOURCES_TIER tiled_tier;
     D3D12_PACKED_MIP_INFO packed_mip_info;
     D3D12_SUBRESOURCE_TILING tilings[10];
     D3D12_TILE_RANGE_FLAGS tile_flags[8];
@@ -36956,6 +36961,13 @@ static void test_update_tile_mappings(void)
     if (!init_test_context(&context, &desc))
         return;

+    if ((tiled_tier = get_tiled_resources_tier(context.device)) < D3D12_TILED_RESOURCES_TIER_1)
+    {
+        skip("Tiled resources not supported by device.\n");
+        destroy_test_context(&context);
+        return;
+    }
+
     descriptor_range.RangeType = D3D12_DESCRIPTOR_RANGE_TYPE_SRV;
     descriptor_range.NumDescriptors = 1;
     descriptor_range.BaseShaderRegister = 0;
@@ -37375,7 +37387,7 @@ static void test_update_tile_mappings(void)
     ID3D12Resource_Release(resource_2);
     ID3D12Resource_Release(array_resource);

-    if (!is_amd_windows_device(context.device) && !is_radv_device(context.device))
+    if (tiled_tier >= D3D12_TILED_RESOURCES_TIER_3)
     {
         /* Test 3D image tile mappings */
         resource_desc.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE3D;
@@ -37567,6 +37579,13 @@ static void test_sparse_buffer_memory_lifetime(void)
     if (!init_compute_test_context(&context))
         return;

+    if (get_tiled_resources_tier(context.device) < D3D12_TILED_RESOURCES_TIER_2)
+    {
+        skip("Tiled resources tier 2 not supported by device.\n");
+        destroy_test_context(&context);
+        return;
+    }
+
     memset(&rs_desc, 0, sizeof(rs_desc));
     memset(root_parameters, 0, sizeof(root_parameters));
     memset(&desc_range, 0, sizeof(desc_range));
@@ -37738,6 +37757,13 @@ static void test_copy_tiles(void)
     if (!init_test_context(&context, &desc))
         return;

+    if (get_tiled_resources_tier(context.device) < D3D12_TILED_RESOURCES_TIER_1)
+    {
+        skip("Tiled resources not supported by device.\n");
+        destroy_test_context(&context);
+        return;
+    }
+
     memset(&heap_desc, 0, sizeof(heap_desc));
     heap_desc.Properties.Type = D3D12_HEAP_TYPE_DEFAULT;
     heap_desc.SizeInBytes = TILE_SIZE * 16;
@@ -37800,7 +37826,7 @@ static void test_copy_tiles(void)
     {
         uint32_t offset = i + 4 * TILE_SIZE / sizeof(*buffer_data);
         set_box(&box, offset, 0, 0, offset + 1, 1, 1);
-        todo check_readback_data_uint(&rb.rb, &box, buffer_data[i + buffer_offset / sizeof(*buffer_data)], 0);
+        check_readback_data_uint(&rb.rb, &box, buffer_data[i + buffer_offset / sizeof(*buffer_data)], 0);
     }

release_resource_readback(&rb);
@@ -37829,7 +37855,7 @@ static void test_copy_tiles(void)
     {
         uint32_t offset = i + (TILE_SIZE + buffer_offset) / sizeof(*buffer_data);
         set_box(&box, i, 0, 0, i + 1, 1, 1);
-        todo check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+        check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
     }

release_resource_readback(&rb);
@@ -37892,7 +37918,7 @@ static void test_copy_tiles(void)
                 uint32_t offset = image_tiles[i].tile_idx * TILE_SIZE / sizeof(*buffer_data) + 128 * y + x;
                 set_box(&box, 128 * image_tiles[i].x + x, 128 * image_tiles[i].y + y, 0,
                         128 * image_tiles[i].x + x + 1, 128 * image_tiles[i].y + y + 1, 1);
-                todo_if(buffer_data[offset]) check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+                check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
             }
         }
     }
@@ -37929,7 +37955,7 @@ static void test_copy_tiles(void)
         {
             uint32_t offset = image_tiles[i].tile_idx * TILE_SIZE / sizeof(uint32_t) + x;
             set_box(&box, offset, 0, 0, offset + 1, 1, 1);
-            todo_if(buffer_data[offset]) check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
+            check_readback_data_uint(&rb.rb, &box, buffer_data[offset], 0);
         }
     }

[PATCH 0/21] MR221: Draft: vkd3d: Implement tiled resources - closed

[Conor McCarthy (＠cmccarthy)]

This merge request was closed by Conor McCarthy.