From: Conor McCarthy cmccarthy@codeweavers.com
--- 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, ®ion_size, &coordinate, ®ion_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(®ion_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); }