vkd3d-shader: Add offset buffer path to DXIL.

Signed-off-by: Hans-Kristian Arntzen <post@arntzen-software.no>

include/private/hashmap.h

@@ -50,14 +50,9 @@ struct hash_map
     uint32_t used_count;
 };
 
-static inline void *hash_map_ptr_offset(void *ptr, size_t offset)
-{
-    return ((char*)ptr) + offset;
-}
-
 static inline struct hash_map_entry *hash_map_get_entry(struct hash_map *hash_map, uint32_t entry_idx)
 {
-    return hash_map_ptr_offset(hash_map->entries, hash_map->entry_size * entry_idx);
+    return void_ptr_offset(hash_map->entries, hash_map->entry_size * entry_idx);
 }
 
 static inline uint32_t hash_map_get_entry_idx(struct hash_map *hash_map, uint32_t hash_value)
@@ -98,7 +93,7 @@ static inline bool hash_map_grow(struct hash_map *hash_map)
     for (i = 0; i < old_count; i++)
     {
         /* Relocate existing entries one by one */
-        struct hash_map_entry *old_entry = hash_map_ptr_offset(old_entries, i * hash_map->entry_size);
+        struct hash_map_entry *old_entry = void_ptr_offset(old_entries, i * hash_map->entry_size);
 
         if (old_entry->flags & HASH_MAP_ENTRY_OCCUPIED)
         {

include/private/vkd3d_common.h

@@ -271,4 +271,9 @@ static inline size_t vkd3d_wcslen(const WCHAR *wstr, size_t wchar_size)
     }
 }
 
+static inline void *void_ptr_offset(void *ptr, size_t offset)
+{
+    return ((char*)ptr) + offset;
+}
+
 #endif  /* __VKD3D_COMMON_H */

include/vkd3d_shader.h

@@ -187,6 +187,7 @@ enum vkd3d_shader_interface_flag
 {
     VKD3D_SHADER_INTERFACE_PUSH_CONSTANTS_AS_UNIFORM_BUFFER = 0x00000001u,
     VKD3D_SHADER_INTERFACE_BINDLESS_CBV_AS_STORAGE_BUFFER   = 0x00000002u,
+    VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER               = 0x00000004u,
 };
 
 struct vkd3d_shader_interface_info
@@ -205,6 +206,8 @@ struct vkd3d_shader_interface_info
 
     /* Ignored unless VKD3D_SHADER_INTERFACE_PUSH_CONSTANTS_AS_UNIFORM_BUFFER is set */
     const struct vkd3d_shader_descriptor_binding *push_constant_ubo_binding;
+    /* Ignored unless VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER is set */
+    const struct vkd3d_shader_descriptor_binding *offset_buffer_binding;
 };
 
 struct vkd3d_shader_transform_feedback_element

libs/vkd3d-shader/dxil.c

@@ -127,21 +127,19 @@ static dxil_spv_bool dxil_srv_remap(void *userdata, const dxil_spv_d3d_binding *
 {
     const struct vkd3d_shader_interface_info *shader_interface_info = userdata;
     unsigned int resource_flags, resource_flags_ssbo;
-    bool use_ssbo;
 
     resource_flags_ssbo = dxil_resource_flags_from_kind(d3d_binding->kind, true);
     resource_flags = dxil_resource_flags_from_kind(d3d_binding->kind, false);
-
-    use_ssbo = resource_flags_ssbo != resource_flags;
+    bool use_ssbo = resource_flags_ssbo != resource_flags;
 
     if (use_ssbo && dxil_remap(shader_interface_info, VKD3D_SHADER_DESCRIPTOR_TYPE_SRV,
             d3d_binding, &vk_binding->buffer_binding, resource_flags_ssbo))
     {
         vk_binding->buffer_binding.descriptor_type = DXIL_SPV_VULKAN_DESCRIPTOR_TYPE_SSBO;
-        if (d3d_binding->alignment < shader_interface_info->min_ssbo_alignment)
+        if (shader_interface_info->flags & VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER)
         {
-            FIXME("Shader declares resource with alignment of %u bytes, but implementation only supports %u.\n",
-                  d3d_binding->alignment, shader_interface_info->min_ssbo_alignment);
+            vk_binding->offset_binding.set = shader_interface_info->offset_buffer_binding->set;
+            vk_binding->offset_binding.binding = shader_interface_info->offset_buffer_binding->binding;
         }
         return DXIL_SPV_TRUE;
     }
@@ -226,15 +224,10 @@ static dxil_spv_bool dxil_uav_remap(void *userdata, const dxil_spv_uav_d3d_bindi
 {
     const struct vkd3d_shader_interface_info *shader_interface_info = userdata;
     unsigned int resource_flags, resource_flags_ssbo;
-    bool use_ssbo;
 
     resource_flags_ssbo = dxil_resource_flags_from_kind(d3d_binding->d3d_binding.kind, true);
     resource_flags = dxil_resource_flags_from_kind(d3d_binding->d3d_binding.kind, false);
-
-    if (resource_flags != resource_flags_ssbo)
-        use_ssbo = d3d_binding->d3d_binding.alignment >= shader_interface_info->min_ssbo_alignment;
-    else
-        use_ssbo = false;
+    bool use_ssbo = resource_flags != resource_flags_ssbo;
 
     if (use_ssbo)
     {
@@ -242,6 +235,11 @@ static dxil_spv_bool dxil_uav_remap(void *userdata, const dxil_spv_uav_d3d_bindi
                 &vk_binding->buffer_binding, resource_flags_ssbo))
         {
             vk_binding->buffer_binding.descriptor_type = DXIL_SPV_VULKAN_DESCRIPTOR_TYPE_SSBO;
+            if (shader_interface_info->flags & VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER)
+            {
+                vk_binding->offset_binding.set = shader_interface_info->offset_buffer_binding->set;
+                vk_binding->offset_binding.binding = shader_interface_info->offset_buffer_binding->binding;
+            }
         }
         else if (!dxil_remap(shader_interface_info, VKD3D_SHADER_DESCRIPTOR_TYPE_UAV, &d3d_binding->d3d_binding,
                 &vk_binding->buffer_binding, resource_flags))

libs/vkd3d-shader/spirv.c

@@ -2195,6 +2195,8 @@ struct vkd3d_dxbc_compiler
     struct vkd3d_shader_global_binding *global_bindings;
     size_t global_bindings_size;
     size_t global_binding_count;
+
+    uint32_t offset_buffer_var_id;
 };
 
 static bool shader_is_sm_5_1(const struct vkd3d_dxbc_compiler *compiler)
@@ -5446,6 +5448,33 @@ static void vkd3d_dxbc_compiler_emit_dcl_indexable_temp(struct vkd3d_dxbc_compil
     vkd3d_dxbc_compiler_put_symbol(compiler, &reg_symbol);
 }
 
+static void vkd3d_dxbc_compiler_emit_offset_buffer(struct vkd3d_dxbc_compiler *compiler)
+{
+    const struct vkd3d_shader_interface_info *shader_interface = &compiler->shader_interface;
+    struct vkd3d_spirv_builder *builder = &compiler->spirv_builder;
+    uint32_t array_id, struct_id, pointer_id, var_id;
+
+    if (!(shader_interface->flags & VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER))
+        return;
+
+    array_id = vkd3d_spirv_build_op_type_runtime_array(builder,
+            vkd3d_spirv_get_type_id(builder, VKD3D_TYPE_UINT, 2));
+    vkd3d_spirv_build_op_decorate1(builder, array_id, SpvDecorationArrayStride, 8);
+
+    struct_id = vkd3d_spirv_build_op_type_struct(builder, &array_id, 1);
+    vkd3d_spirv_build_op_decorate(builder, struct_id, SpvDecorationBufferBlock, NULL, 0);
+    vkd3d_spirv_build_op_member_decorate1(builder, struct_id, 0, SpvDecorationOffset, 0);
+    vkd3d_spirv_build_op_name(builder, struct_id, "offset_buf");
+
+    pointer_id = vkd3d_spirv_build_op_type_pointer(builder, SpvStorageClassUniform, struct_id);
+    var_id = vkd3d_spirv_build_op_variable(builder, &builder->global_stream, pointer_id, SpvStorageClassUniform, 0);
+    vkd3d_spirv_build_op_decorate1(builder, var_id, SpvDecorationDescriptorSet, shader_interface->offset_buffer_binding->set);
+    vkd3d_spirv_build_op_decorate1(builder, var_id, SpvDecorationBinding, shader_interface->offset_buffer_binding->binding);
+    vkd3d_spirv_build_op_decorate(builder, var_id, SpvDecorationNonWritable, NULL, 0);
+
+    compiler->offset_buffer_var_id = var_id;
+}
+
 static void vkd3d_dxbc_compiler_emit_push_constant_buffers(struct vkd3d_dxbc_compiler *compiler)
 {
     uint32_t uint_id, float_id, struct_id, pointer_type_id, var_id;
@@ -8416,6 +8445,77 @@ static uint32_t vkd3d_dxbc_compiler_emit_raw_structured_addressing(
         return offset_id;
 }
 
+static uint32_t vkd3d_dxbc_compiler_get_ssbo_bounds(struct vkd3d_dxbc_compiler *compiler,
+        const struct vkd3d_shader_register *reg, const struct vkd3d_shader_resource_binding *binding)
+{
+    struct vkd3d_spirv_builder *builder = &compiler->spirv_builder;
+    uint32_t bounds_id, vec2_ptr_id, vec2_type_id;
+    uint32_t indices[2];
+
+    vec2_type_id = vkd3d_spirv_get_type_id(builder, VKD3D_TYPE_UINT, 2);
+    vec2_ptr_id = vkd3d_spirv_get_op_type_pointer(builder, SpvStorageClassUniform, vec2_type_id);
+
+    indices[0] = vkd3d_dxbc_compiler_get_constant_uint(compiler, 0);
+    indices[1] = vkd3d_dxbc_compiler_get_resource_index(compiler, reg, binding);
+
+    /* returns (offset, length) in bytes */
+    bounds_id = vkd3d_spirv_build_op_load(builder, vec2_type_id,
+            vkd3d_spirv_build_op_access_chain(builder, vec2_ptr_id,
+                    compiler->offset_buffer_var_id, indices, ARRAY_SIZE(indices)),
+            SpvMemoryAccessMaskNone);
+
+    return bounds_id;
+}
+
+static uint32_t vkd3d_dxbc_compiler_adjust_ssbo_offset(struct vkd3d_dxbc_compiler *compiler,
+        const struct vkd3d_shader_register *reg, uint32_t coordinate_id)
+{
+    struct vkd3d_spirv_builder *builder = &compiler->spirv_builder;
+    uint32_t shift_id, offset_id, length_id, bounds_id, cond_id;
+    uint32_t uint_type_id, bool_type_id;
+    const struct vkd3d_symbol *symbol;
+    unsigned int alignment;
+
+    if (!(compiler->shader_interface.flags & VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER))
+        return coordinate_id;
+
+    symbol = vkd3d_dxbc_compiler_find_resource(compiler, reg);
+
+    if (symbol->info.resource.raw)
+        alignment = 16;
+    else
+        alignment = 4 * (symbol->info.resource.structure_stride & -symbol->info.resource.structure_stride);
+
+    /* Assume that offset is 0 and size matches the descriptor size */
+    if (alignment >= compiler->shader_interface.min_ssbo_alignment)
+        return coordinate_id;
+
+    bool_type_id = vkd3d_spirv_get_type_id(builder, VKD3D_TYPE_BOOL, 1);
+    uint_type_id = vkd3d_spirv_get_type_id(builder, VKD3D_TYPE_UINT, 1);
+
+    bounds_id = vkd3d_dxbc_compiler_get_ssbo_bounds(compiler, reg, symbol->info.resource.resource_binding);
+
+    shift_id = vkd3d_dxbc_compiler_get_constant_uint(compiler, 2);
+    offset_id = vkd3d_spirv_build_op_shift_right_logical(builder, uint_type_id,
+            vkd3d_spirv_build_op_composite_extract1(builder, uint_type_id, bounds_id, 0), shift_id);
+    length_id = vkd3d_spirv_build_op_shift_right_logical(builder, uint_type_id,
+            vkd3d_spirv_build_op_composite_extract1(builder, uint_type_id, bounds_id, 1), shift_id);
+
+    /* cond = offset < length */
+    cond_id = vkd3d_spirv_build_op_uless_than(builder, bool_type_id, coordinate_id, length_id);
+
+    /* In case of out-of-bounds access, set offset to a number that we
+     * expect to be out-of-bounds of the actual Vulkan resource as well.
+     * 0x3ffffffc is the largest offset value we can safely use without
+     * overflowing 32-bit address space, since this is a DWORD offset
+     * and we may access a total of 16 bytes starting at that offset. */
+    coordinate_id = vkd3d_spirv_build_op_select(builder, uint_type_id, cond_id,
+            vkd3d_spirv_build_op_iadd(builder, uint_type_id, coordinate_id, offset_id),
+            vkd3d_dxbc_compiler_get_constant_uint(compiler, 0x3ffffffc));
+
+    return coordinate_id;
+}
+
 static void vkd3d_dxbc_compiler_emit_ld_raw_structured_srv_uav(struct vkd3d_dxbc_compiler *compiler,
         const struct vkd3d_shader_instruction *instruction)
 {
@@ -8462,6 +8562,9 @@ static void vkd3d_dxbc_compiler_emit_ld_raw_structured_srv_uav(struct vkd3d_dxbc
     base_coordinate_id = vkd3d_dxbc_compiler_emit_raw_structured_addressing(compiler,
             type_id, image.structure_stride, &src[0], VKD3DSP_WRITEMASK_0, &src[1], VKD3DSP_WRITEMASK_0);
 
+    if (image.ssbo)
+        base_coordinate_id = vkd3d_dxbc_compiler_adjust_ssbo_offset(compiler, &resource->reg, base_coordinate_id);
+
     texel_type_id = vkd3d_spirv_get_type_id(builder, image.sampled_type, VKD3D_VEC4_SIZE);
     result_type_id = is_sparse_op ? vkd3d_spirv_get_sparse_result_type(builder, texel_type_id) : texel_type_id;
     assert(dst->write_mask & VKD3DSP_WRITEMASK_ALL);
@@ -8594,6 +8697,9 @@ static void vkd3d_dxbc_compiler_emit_store_uav_raw_structured(struct vkd3d_dxbc_
     base_coordinate_id = vkd3d_dxbc_compiler_emit_raw_structured_addressing(compiler,
             type_id, image.structure_stride, &src[0], VKD3DSP_WRITEMASK_0, &src[1], VKD3DSP_WRITEMASK_0);
 
+    if (image.ssbo)
+        base_coordinate_id = vkd3d_dxbc_compiler_adjust_ssbo_offset(compiler, &dst->reg, base_coordinate_id);
+
     texel = &src[instruction->src_count - 1];
     assert(texel->reg.data_type == VKD3D_DATA_UINT);
     val_id = vkd3d_dxbc_compiler_emit_load_src(compiler, texel, dst->write_mask);
@@ -8936,6 +9042,9 @@ static void vkd3d_dxbc_compiler_emit_atomic_instruction(struct vkd3d_dxbc_compil
         coordinate_id = vkd3d_dxbc_compiler_emit_raw_structured_addressing(compiler,
                 type_id, structure_stride, &src[0], VKD3DSP_WRITEMASK_0,
                 &src[0], VKD3DSP_WRITEMASK_1);
+
+        if (resource->reg.type != VKD3DSPR_GROUPSHAREDMEM && image.ssbo)
+            coordinate_id = vkd3d_dxbc_compiler_adjust_ssbo_offset(compiler, &resource->reg, coordinate_id);
     }
     else
     {
@@ -8999,7 +9108,7 @@ static void vkd3d_dxbc_compiler_emit_bufinfo(struct vkd3d_dxbc_compiler *compile
     struct vkd3d_spirv_builder *builder = &compiler->spirv_builder;
     const struct vkd3d_shader_dst_param *dst = instruction->dst;
     const struct vkd3d_shader_src_param *src = instruction->src;
-    uint32_t type_id, val_id, stride_id;
+    uint32_t type_id, val_id, stride_id, bounds_id;
     struct vkd3d_shader_image image;
     uint32_t constituents[2];
     unsigned int write_mask;
@@ -9012,10 +9121,22 @@ static void vkd3d_dxbc_compiler_emit_bufinfo(struct vkd3d_dxbc_compiler *compile
 
     if (image.ssbo)
     {
-        if (src->reg.modifier == VKD3DSPRM_NONUNIFORM)
-            vkd3d_dxbc_compiler_decorate_nonuniform(compiler, image.id);
+        if (compiler->shader_interface.flags & VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER)
+        {
+            const struct vkd3d_symbol *symbol = vkd3d_dxbc_compiler_find_resource(compiler, &src->reg);
+            bounds_id = vkd3d_dxbc_compiler_get_ssbo_bounds(compiler, &src->reg, symbol->info.resource.resource_binding);
 
-        val_id = vkd3d_spirv_build_op_array_length(builder, type_id, image.id, 0);
+            val_id = vkd3d_spirv_build_op_shift_right_logical(builder, type_id,
+                    vkd3d_spirv_build_op_composite_extract1(builder, type_id, bounds_id, 1),
+                    vkd3d_dxbc_compiler_get_constant_uint(compiler, 2));
+        }
+        else
+        {
+            if (src->reg.modifier == VKD3DSPRM_NONUNIFORM)
+                vkd3d_dxbc_compiler_decorate_nonuniform(compiler, image.id);
+
+            val_id = vkd3d_spirv_build_op_array_length(builder, type_id, image.id, 0);
+        }
     }
     else
         val_id = vkd3d_spirv_build_op_image_query_size(builder, type_id, image.image_id);
@@ -9457,6 +9578,7 @@ static void vkd3d_dxbc_compiler_emit_double_conversion(struct vkd3d_dxbc_compile
 /* This function is called after declarations are processed. */
 static void vkd3d_dxbc_compiler_emit_main_prolog(struct vkd3d_dxbc_compiler *compiler)
 {
+    vkd3d_dxbc_compiler_emit_offset_buffer(compiler);
     vkd3d_dxbc_compiler_emit_push_constant_buffers(compiler);
 }
 

libs/vkd3d/command.c

@@ -5164,7 +5164,7 @@ static void d3d12_command_list_set_root_descriptor(struct d3d12_command_list *li
     bool null_descriptors, ssbo;
     VkDeviceSize max_range;
 
-    ssbo = d3d12_device_use_ssbo_raw_buffer(list->device);
+    ssbo = d3d12_device_use_ssbo_root_descriptors(list->device);
     root_parameter = root_signature_get_root_descriptor(root_signature, index);
     descriptor = &bindings->root_descriptors[root_parameter->descriptor.packed_descriptor];
     null_descriptors = list->device->device_info.robustness2_features.nullDescriptor;
@@ -5711,7 +5711,7 @@ struct vkd3d_clear_uav_info
     } u;
 };
 
-static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
+static void d3d12_command_list_clear_uav(struct d3d12_command_list *list, const struct d3d12_desc *desc,
         struct d3d12_resource *resource, const struct vkd3d_clear_uav_info *args,
         const VkClearColorValue *clear_color, UINT rect_count, const D3D12_RECT *rects)
 {
@@ -5724,6 +5724,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
     D3D12_RECT full_rect, curr_rect;
     VkWriteDescriptorSet write_set;
     VkExtent3D workgroup_size;
+    uint32_t extra_offset;
 
     d3d12_command_list_track_resource_usage(list, resource);
     d3d12_command_list_end_current_render_pass(list, false);
@@ -5802,6 +5803,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
     full_rect.right = d3d12_resource_desc_get_width(&resource->desc, miplevel_idx);
     full_rect.top = 0;
     full_rect.bottom = d3d12_resource_desc_get_height(&resource->desc, miplevel_idx);
+    extra_offset = 0;
 
     if (d3d12_resource_is_buffer(resource))
     {
@@ -5812,6 +5814,12 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
                     : sizeof(uint32_t);  /* structured buffer */
             full_rect.right = args->u.view->info.buffer.size / byte_count;
         }
+        else if (list->device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+        {
+            const struct vkd3d_bound_ssbo_range *ranges = desc->heap->ssbo_ranges.host_ptr;
+            extra_offset = ranges[desc->heap_offset].offset / sizeof(uint32_t);
+            full_rect.right = ranges[desc->heap_offset].length / sizeof(uint32_t);
+        }
         else
             full_rect.right = args->u.buffer.range / sizeof(uint32_t);
     }
@@ -5840,7 +5848,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
                 continue;
         }
 
-        clear_args.offset.x = curr_rect.left;
+        clear_args.offset.x = curr_rect.left + extra_offset;
         clear_args.offset.y = curr_rect.top;
         clear_args.extent.width = curr_rect.right - curr_rect.left;
         clear_args.extent.height = curr_rect.bottom - curr_rect.top;
@@ -5956,7 +5964,7 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(d3
         }
     }
 
-    d3d12_command_list_clear_uav(list, resource_impl, &args, &color, rect_count, rects);
+    d3d12_command_list_clear_uav(list, desc, resource_impl, &args, &color, rect_count, rects);
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(d3d12_command_list_iface *iface,
@@ -5977,7 +5985,7 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(d
     resource_impl = unsafe_impl_from_ID3D12Resource(resource);
 
     vkd3d_clear_uav_info_from_desc(&args, desc);
-    d3d12_command_list_clear_uav(list, resource_impl, &args, &color, rect_count, rects);
+    d3d12_command_list_clear_uav(list, desc, resource_impl, &args, &color, rect_count, rects);
 }
 
 static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(d3d12_command_list_iface *iface,

libs/vkd3d/resource.c

@@ -3478,9 +3478,10 @@ void d3d12_desc_copy(struct d3d12_desc *dst, struct d3d12_desc *src,
 
         if (metadata.flags & VKD3D_DESCRIPTOR_FLAG_UAV_COUNTER)
         {
-            if (dst->heap->uav_counters.data)
+            if (dst->heap->uav_counters.host_ptr)
             {
-                dst->heap->uav_counters.data[dst->heap_offset] = src->counter_address;
+                VkDeviceAddress *counter_addresses = dst->heap->uav_counters.host_ptr;
+                counter_addresses[dst->heap_offset] = src->counter_address;
                 dst->counter_address = src->counter_address;
             }
             else
@@ -3501,6 +3502,13 @@ void d3d12_desc_copy(struct d3d12_desc *dst, struct d3d12_desc *src,
             }
         }
 
+        if (metadata.flags & VKD3D_DESCRIPTOR_FLAG_SSBO_OFFSET)
+        {
+            const struct vkd3d_bound_ssbo_range *src_ssbo_ranges = src->heap->ssbo_ranges.host_ptr;
+            struct vkd3d_bound_ssbo_range *dst_ssbo_ranges = dst->heap->ssbo_ranges.host_ptr;
+            dst_ssbo_ranges[dst->heap_offset] = src_ssbo_ranges[src->heap_offset];
+        }
+
         if (copy_count)
             VK_CALL(vkUpdateDescriptorSets(device->vk_device, 0, NULL, copy_count, vk_copies));
     }
@@ -3947,6 +3955,43 @@ static bool vkd3d_buffer_srv_use_raw_ssbo(struct d3d12_device *device,
             ((desc->Format == DXGI_FORMAT_UNKNOWN && desc->Buffer.StructureByteStride) || raw);
 }
 
+static void vkd3d_buffer_view_get_bound_range(struct d3d12_desc *descriptor,
+        struct d3d12_device *device, struct d3d12_resource *resource,
+        VkDeviceSize offset, VkDeviceSize range, VkDescriptorBufferInfo *vk_buffer)
+{
+    struct vkd3d_bound_ssbo_range ssbo_range;
+
+    if (resource)
+    {
+        VkDeviceSize alignment = d3d12_device_get_ssbo_alignment(device);
+        VkDeviceSize aligned_begin = offset & ~(alignment - 1);
+        VkDeviceSize aligned_end = min((offset + range + alignment - 1) & ~(alignment - 1), resource->desc.Width);
+
+        /* heap_offset is guaranteed to have 64KiB alignment */
+        vk_buffer->buffer = resource->vk_buffer;
+        vk_buffer->offset = resource->heap_offset + aligned_begin;
+        vk_buffer->range = aligned_end - aligned_begin;
+
+        ssbo_range.offset = offset - aligned_begin;
+        ssbo_range.length = range;
+    }
+    else
+    {
+        vk_buffer->buffer = VK_NULL_HANDLE;
+        vk_buffer->offset = 0;
+        vk_buffer->range = 0;
+
+        ssbo_range.offset = 0;
+        ssbo_range.length = 0;
+    }
+
+    if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+    {
+        struct vkd3d_bound_ssbo_range *ssbo_ranges = descriptor->heap->ssbo_ranges.host_ptr;
+        ssbo_ranges[descriptor->heap_offset] = ssbo_range;
+    }
+}
+
 static void vkd3d_create_buffer_srv(struct d3d12_desc *descriptor,
         struct d3d12_device *device, struct d3d12_resource *resource,
         const D3D12_SHADER_RESOURCE_VIEW_DESC *desc)
@@ -3973,26 +4018,12 @@ static void vkd3d_create_buffer_srv(struct d3d12_desc *descriptor,
 
     if (vkd3d_buffer_srv_use_raw_ssbo(device, desc))
     {
-        if (resource)
-        {
-            uint32_t stride = desc->Format == DXGI_FORMAT_UNKNOWN
-                    ? desc->Buffer.StructureByteStride : sizeof(uint32_t);
-            descriptor_info.buffer.buffer = resource->vk_buffer;
-            descriptor_info.buffer.offset = desc->Buffer.FirstElement * stride + resource->heap_offset;
-            descriptor_info.buffer.range = desc->Buffer.NumElements * stride;
+        VkDeviceSize stride = desc->Format == DXGI_FORMAT_UNKNOWN
+                ? desc->Buffer.StructureByteStride : sizeof(uint32_t);
 
-            if (descriptor_info.buffer.offset & (d3d12_device_get_ssbo_alignment(device) - 1))
-            {
-                FIXME("Emitting SSBO at offset #%"PRIx64", but needs alignment of %"PRIu64" bytes.\n",
-                      descriptor_info.buffer.offset, d3d12_device_get_ssbo_alignment(device));
-            }
-        }
-        else
-        {
-            descriptor_info.buffer.buffer = VK_NULL_HANDLE;
-            descriptor_info.buffer.offset = 0;
-            descriptor_info.buffer.range = 0;
-        }
+        vkd3d_buffer_view_get_bound_range(descriptor, device, resource,
+                desc->Buffer.FirstElement * stride, desc->Buffer.NumElements * stride,
+                &descriptor_info.buffer);
 
         descriptor->info.buffer = descriptor_info.buffer;
         descriptor->metadata.cookie = resource ? resource->cookie : 0;
@@ -4000,6 +4031,9 @@ static void vkd3d_create_buffer_srv(struct d3d12_desc *descriptor,
                 VKD3D_BINDLESS_SET_SRV | VKD3D_BINDLESS_SET_RAW_SSBO);
         descriptor->metadata.flags = VKD3D_DESCRIPTOR_FLAG_DEFINED;
 
+        if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+            descriptor->metadata.flags |= VKD3D_DESCRIPTOR_FLAG_SSBO_OFFSET;
+
         vk_descriptor_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
     }
     else
@@ -4286,33 +4320,21 @@ static void vkd3d_create_buffer_uav(struct d3d12_desc *descriptor, struct d3d12_
     if (vkd3d_buffer_uav_use_raw_ssbo(device, desc))
     {
         VkDescriptorBufferInfo *buffer_info = &descriptor_info[vk_write_count].buffer;
+        VkDeviceSize stride = desc->Format == DXGI_FORMAT_UNKNOWN
+                ? desc->Buffer.StructureByteStride : sizeof(uint32_t);
 
-        if (resource)
-        {
-            uint32_t stride = desc->Format == DXGI_FORMAT_UNKNOWN
-                    ? desc->Buffer.StructureByteStride : sizeof(uint32_t);
-            buffer_info->buffer = resource->vk_buffer;
-            buffer_info->offset = desc->Buffer.FirstElement * stride + resource->heap_offset;
-            buffer_info->range = desc->Buffer.NumElements * stride;
-
-            if (buffer_info->offset & (d3d12_device_get_ssbo_alignment(device) - 1))
-            {
-                FIXME("Emitting SSBO at offset #%"PRIx64", but needs alignment of %"PRIu64" bytes.\n",
-                      buffer_info->offset, d3d12_device_get_ssbo_alignment(device));
-            }
-        }
-        else
-        {
-            buffer_info->buffer = VK_NULL_HANDLE;
-            buffer_info->offset = 0;
-            buffer_info->range = 0;
-        }
+        vkd3d_buffer_view_get_bound_range(descriptor, device, resource,
+                desc->Buffer.FirstElement * stride, desc->Buffer.NumElements * stride,
+                buffer_info);
 
         descriptor->info.buffer = *buffer_info;
         descriptor->metadata.cookie = resource ? resource->cookie : 0;
         descriptor->metadata.binding = vkd3d_bindless_state_find_set(&device->bindless_state,
                 VKD3D_BINDLESS_SET_UAV | VKD3D_BINDLESS_SET_RAW_SSBO);
-        descriptor->metadata.flags = VKD3D_DESCRIPTOR_FLAG_DEFINED;
+        descriptor->metadata.flags = VKD3D_DESCRIPTOR_FLAG_DEFINED | VKD3D_DESCRIPTOR_FLAG_UAV_COUNTER;
+
+        if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+            descriptor->metadata.flags |= VKD3D_DESCRIPTOR_FLAG_SSBO_OFFSET;
 
         vk_descriptor_type = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER;
     }
@@ -4382,8 +4404,9 @@ static void vkd3d_create_buffer_uav(struct d3d12_desc *descriptor, struct d3d12_
 
     if (device->bindless_state.flags & VKD3D_RAW_VA_UAV_COUNTER)
     {
+        VkDeviceAddress *counter_addresses = descriptor->heap->uav_counters.host_ptr;
         uint32_t descriptor_index = d3d12_desc_heap_offset(descriptor);
-        descriptor->heap->uav_counters.data[descriptor_index] = uav_counter_address;
+        counter_addresses[descriptor_index] = uav_counter_address;
         descriptor->counter_address = uav_counter_address;
     }
     else
@@ -5400,48 +5423,98 @@ static HRESULT d3d12_descriptor_heap_create_descriptor_set(struct d3d12_descript
     return S_OK;
 }
 
-static HRESULT d3d12_descriptor_heap_create_uav_counter_buffer(struct d3d12_descriptor_heap *descriptor_heap,
-        struct d3d12_descriptor_heap_uav_counters *uav_counters)
+static void d3d12_descriptor_heap_get_buffer_range(struct d3d12_descriptor_heap *descriptor_heap,
+        VkDeviceSize *offset, VkDeviceSize size, struct vkd3d_host_visible_buffer_range *range)
+{
+    if (size)
+    {
+        range->descriptor.buffer = descriptor_heap->vk_buffer;
+        range->descriptor.offset = *offset;
+        range->descriptor.range = size;
+        range->host_ptr = void_ptr_offset(descriptor_heap->host_memory, *offset);
+
+        *offset += size;
+    }
+    else
+    {
+        range->descriptor.buffer = VK_NULL_HANDLE;
+        range->descriptor.offset = 0;
+        range->descriptor.range = VK_WHOLE_SIZE;
+        range->host_ptr = NULL;
+    }
+}
+
+static HRESULT d3d12_descriptor_heap_init_data_buffer(struct d3d12_descriptor_heap *descriptor_heap,
+        struct d3d12_device *device, const D3D12_DESCRIPTOR_HEAP_DESC *desc)
 {
     const struct vkd3d_vk_device_procs *vk_procs = &descriptor_heap->device->vk_procs;
-    struct d3d12_device *device = descriptor_heap->device;
+    VkDeviceSize alignment = max(device->device_info.properties2.properties.limits.minStorageBufferOffsetAlignment,
+            device->device_info.properties2.properties.limits.nonCoherentAtomSize);
+    VkDeviceSize uav_counter_size = 0, offset_buffer_size = 0;
+    VkDeviceSize buffer_size, offset;
     D3D12_HEAP_PROPERTIES heap_info;
     D3D12_RESOURCE_DESC buffer_desc;
     D3D12_HEAP_FLAGS heap_flags;
     VkResult vr;
     HRESULT hr;
 
-    /* concurrently accessible storage buffer */
-    memset(&buffer_desc, 0, sizeof(buffer_desc));
-    buffer_desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
-    buffer_desc.Width = descriptor_heap->desc.NumDescriptors * sizeof(VkDeviceAddress);
-    buffer_desc.Height = 1;
-    buffer_desc.DepthOrArraySize = 1;
-    buffer_desc.MipLevels = 1;
-    buffer_desc.SampleDesc.Count = 1;
-    buffer_desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
-    buffer_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
-
-    /* host-visible device memory */
-    memset(&heap_info, 0, sizeof(heap_info));
-    heap_info.Type = D3D12_HEAP_TYPE_UPLOAD;
-
-    heap_flags = D3D12_HEAP_FLAG_ALLOW_ONLY_BUFFERS;
-
-    if (FAILED(hr = vkd3d_create_buffer(device, &heap_info, heap_flags, &buffer_desc, &uav_counters->vk_buffer)))
-        return hr;
-
-    if (FAILED(hr = vkd3d_allocate_buffer_memory(device, uav_counters->vk_buffer, NULL,
-            &heap_info, heap_flags, &uav_counters->vk_memory, NULL, NULL)))
-        return hr;
-
-    if ((vr = VK_CALL(vkMapMemory(device->vk_device, uav_counters->vk_memory,
-            0, VK_WHOLE_SIZE, 0, (void **)&uav_counters->data))))
+    if (desc->Type == D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV)
     {
-        ERR("Failed to map UAV counter address buffer, vr %d.\n", vr);
-        return hresult_from_vk_result(vr);
+        if (device->bindless_state.flags & VKD3D_RAW_VA_UAV_COUNTER)
+            uav_counter_size = align(desc->NumDescriptors * sizeof(VkDeviceAddress), alignment);
+
+        if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+            offset_buffer_size = align(desc->NumDescriptors * sizeof(struct vkd3d_bound_ssbo_range), alignment);
     }
 
+    buffer_size = uav_counter_size + offset_buffer_size;
+
+    if (!buffer_size)
+        return S_OK;
+
+    if (desc->Flags & D3D12_DESCRIPTOR_HEAP_FLAG_SHADER_VISIBLE)
+    {
+        memset(&buffer_desc, 0, sizeof(buffer_desc));
+        buffer_desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER;
+        buffer_desc.Width = buffer_size;
+        buffer_desc.Height = 1;
+        buffer_desc.DepthOrArraySize = 1;
+        buffer_desc.MipLevels = 1;
+        buffer_desc.SampleDesc.Count = 1;
+        buffer_desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR;
+        buffer_desc.Flags = D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS;
+
+        /* host-visible device memory */
+        memset(&heap_info, 0, sizeof(heap_info));
+        heap_info.Type = D3D12_HEAP_TYPE_UPLOAD;
+
+        heap_flags = D3D12_HEAP_FLAG_ALLOW_ONLY_BUFFERS;
+
+        if (FAILED(hr = vkd3d_create_buffer(device, &heap_info, heap_flags, &buffer_desc, &descriptor_heap->vk_buffer)))
+            return hr;
+
+        if (FAILED(hr = vkd3d_allocate_buffer_memory(device, descriptor_heap->vk_buffer, NULL,
+                &heap_info, heap_flags, &descriptor_heap->vk_memory, NULL, NULL)))
+            return hr;
+
+        if ((vr = VK_CALL(vkMapMemory(device->vk_device, descriptor_heap->vk_memory,
+                0, VK_WHOLE_SIZE, 0, &descriptor_heap->host_memory))))
+        {
+            ERR("Failed to map buffer, vr %d.\n", vr);
+            return hresult_from_vk_result(vr);
+        }
+    }
+    else
+    {
+        descriptor_heap->vk_memory = VK_NULL_HANDLE;
+        descriptor_heap->vk_buffer = VK_NULL_HANDLE;
+        descriptor_heap->host_memory = vkd3d_calloc(1, buffer_size);
+    }
+
+    offset = 0;
+
+    d3d12_descriptor_heap_get_buffer_range(descriptor_heap, &offset, uav_counter_size, &descriptor_heap->uav_counters);
+    d3d12_descriptor_heap_get_buffer_range(descriptor_heap, &offset, offset_buffer_size, &descriptor_heap->ssbo_ranges);
     return S_OK;
 }
 
@@ -5484,9 +5557,11 @@ static void d3d12_descriptor_heap_update_extra_bindings(struct d3d12_descriptor_
             switch (flag)
             {
                 case VKD3D_BINDLESS_SET_EXTRA_UAV_COUNTER_BUFFER:
-                    vk_buffer->buffer = descriptor_heap->uav_counters.vk_buffer;
-                    vk_buffer->offset = 0;
-                    vk_buffer->range = VK_WHOLE_SIZE;
+                    *vk_buffer = descriptor_heap->uav_counters.descriptor;
+                    break;
+
+                case VKD3D_BINDLESS_SET_EXTRA_SSBO_OFFSET_BUFFER:
+                    *vk_buffer = descriptor_heap->ssbo_ranges.descriptor;
                     break;
 
                 default:
@@ -5537,21 +5612,8 @@ static HRESULT d3d12_descriptor_heap_init(struct d3d12_descriptor_heap *descript
         }
     }
 
-    if (desc->Type == D3D12_DESCRIPTOR_HEAP_TYPE_CBV_SRV_UAV &&
-            (device->bindless_state.flags & VKD3D_RAW_VA_UAV_COUNTER))
-    {
-        if (desc->Flags & D3D12_DESCRIPTOR_HEAP_FLAG_SHADER_VISIBLE)
-        {
-            if (FAILED(hr = d3d12_descriptor_heap_create_uav_counter_buffer(descriptor_heap,
-                    &descriptor_heap->uav_counters)))
-                goto fail;
-        }
-        else if (!(descriptor_heap->uav_counters.data = vkd3d_calloc(desc->NumDescriptors, sizeof(VkDeviceSize))))
-        {
-            ERR("Failed to allocate UAV counter address buffer.\n");
-            goto fail;
-        }
-    }
+    if (FAILED(hr = d3d12_descriptor_heap_init_data_buffer(descriptor_heap, device, desc)))
+        goto fail;
 
     if (desc->Flags & D3D12_DESCRIPTOR_HEAP_FLAG_SHADER_VISIBLE)
         d3d12_descriptor_heap_update_extra_bindings(descriptor_heap, device);
@@ -5648,11 +5710,11 @@ void d3d12_descriptor_heap_cleanup(struct d3d12_descriptor_heap *descriptor_heap
     const struct vkd3d_vk_device_procs *vk_procs = &descriptor_heap->device->vk_procs;
     const struct d3d12_device *device = descriptor_heap->device;
 
-    if (!descriptor_heap->uav_counters.vk_memory)
-        vkd3d_free(descriptor_heap->uav_counters.data);
+    if (!descriptor_heap->vk_memory)
+        vkd3d_free(descriptor_heap->host_memory);
 
-    VK_CALL(vkDestroyBuffer(device->vk_device, descriptor_heap->uav_counters.vk_buffer, NULL));
-    VK_CALL(vkFreeMemory(device->vk_device, descriptor_heap->uav_counters.vk_memory, NULL));
+    VK_CALL(vkDestroyBuffer(device->vk_device, descriptor_heap->vk_buffer, NULL));
+    VK_CALL(vkFreeMemory(device->vk_device, descriptor_heap->vk_memory, NULL));
 
     VK_CALL(vkDestroyDescriptorPool(device->vk_device, descriptor_heap->vk_descriptor_pool, NULL));
 }

libs/vkd3d/state.c

@@ -213,7 +213,7 @@ static enum vkd3d_shader_visibility vkd3d_shader_visibility_from_d3d12(D3D12_SHA
 
 static VkDescriptorType vk_descriptor_type_from_d3d12_root_parameter(struct d3d12_device *device, D3D12_ROOT_PARAMETER_TYPE type)
 {
-    bool use_ssbo = d3d12_device_use_ssbo_raw_buffer(device);
+    bool use_ssbo = d3d12_device_use_ssbo_root_descriptors(device);
 
     switch (type)
     {
@@ -330,6 +330,7 @@ struct d3d12_root_signature_info
     uint32_t push_descriptor_count;
     uint32_t root_constant_count;
     bool has_raw_va_uav_counters;
+    bool has_ssbo_offset_buffer;
 
     uint32_t cost;
 };
@@ -345,6 +346,9 @@ static HRESULT d3d12_root_signature_info_count_descriptors(struct d3d12_root_sig
 
             if (device->bindless_state.flags & VKD3D_BINDLESS_RAW_SSBO)
                 info->binding_count += 1;
+
+            if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+                info->has_ssbo_offset_buffer = true;
             break;
         case D3D12_DESCRIPTOR_RANGE_TYPE_UAV:
             /* separate image + buffer descriptors */
@@ -355,6 +359,9 @@ static HRESULT d3d12_root_signature_info_count_descriptors(struct d3d12_root_sig
 
             if (device->bindless_state.flags & VKD3D_RAW_VA_UAV_COUNTER)
                 info->has_raw_va_uav_counters = true;
+
+            if (device->bindless_state.flags & VKD3D_SSBO_OFFSET_BUFFER)
+                info->has_ssbo_offset_buffer = true;
             break;
         case D3D12_DESCRIPTOR_RANGE_TYPE_CBV:
         case D3D12_DESCRIPTOR_RANGE_TYPE_SAMPLER:
@@ -845,6 +852,15 @@ static HRESULT d3d12_root_signature_init(struct d3d12_root_signature *root_signa
                 &root_signature->uav_counter_binding);
     }
 
+    if (info.has_ssbo_offset_buffer)
+    {
+        root_signature->flags |= VKD3D_ROOT_SIGNATURE_USE_SSBO_OFFSET_BUFFER;
+
+        vkd3d_bindless_state_find_binding(&device->bindless_state,
+                VKD3D_BINDLESS_SET_EXTRA_SSBO_OFFSET_BUFFER,
+                &root_signature->offset_buffer_binding);
+    }
+
     if (FAILED(hr = d3d12_root_signature_init_root_descriptors(root_signature, desc,
                 &info, &root_signature->push_constant_range, &context,
                 &root_signature->vk_root_descriptor_layout)))
@@ -933,6 +949,9 @@ static unsigned int d3d12_root_signature_get_shader_interface_flags(const struct
     if (root_signature->flags & VKD3D_ROOT_SIGNATURE_USE_INLINE_UNIFORM_BLOCK)
         flags |= VKD3D_SHADER_INTERFACE_PUSH_CONSTANTS_AS_UNIFORM_BUFFER;
 
+    if (root_signature->flags & VKD3D_ROOT_SIGNATURE_USE_SSBO_OFFSET_BUFFER)
+        flags |= VKD3D_SHADER_INTERFACE_SSBO_OFFSET_BUFFER;
+
     if (root_signature->device->bindless_state.flags & VKD3D_BINDLESS_CBV_AS_SSBO)
         flags |= VKD3D_SHADER_INTERFACE_BINDLESS_CBV_AS_STORAGE_BUFFER;
 
@@ -1694,6 +1713,7 @@ static HRESULT d3d12_pipeline_state_init_compute(struct d3d12_pipeline_state *st
     shader_interface.push_constant_buffers = root_signature->root_constants;
     shader_interface.push_constant_buffer_count = root_signature->root_constant_count;
     shader_interface.push_constant_ubo_binding = &root_signature->push_constant_ubo_binding;
+    shader_interface.offset_buffer_binding = &root_signature->offset_buffer_binding;
 
     if ((hr = vkd3d_create_pipeline_cache_from_d3d12_desc(device, &desc->cached_pso, &state->vk_pso_cache)) < 0)
     {
@@ -2527,6 +2547,7 @@ static HRESULT d3d12_pipeline_state_init_graphics(struct d3d12_pipeline_state *s
     shader_interface.push_constant_buffers = root_signature->root_constants;
     shader_interface.push_constant_buffer_count = root_signature->root_constant_count;
     shader_interface.push_constant_ubo_binding = &root_signature->push_constant_ubo_binding;
+    shader_interface.offset_buffer_binding = &root_signature->offset_buffer_binding;
 
     graphics->patch_vertex_count = 0;
 
@@ -3436,13 +3457,17 @@ static uint32_t vkd3d_bindless_state_get_bindless_flags(struct d3d12_device *dev
 
     /* Normally, we would be able to use SSBOs conditionally even when maxSSBOAlignment > 4, but
      * applications (RE2 being one example) are of course buggy and don't match descriptor and shader usage of resources,
-     * so we cannot rely on alignment analysis to select the appropriate resource type.
-     * TODO: Implement an offset buffer system so that we can remove the minStorageBufferOffsetAlignment requirement. */
+     * so we cannot rely on alignment analysis to select the appropriate resource type. */
     if (device_info->descriptor_indexing_properties.maxPerStageDescriptorUpdateAfterBindStorageBuffers >= 1000000 &&
             device_info->descriptor_indexing_features.descriptorBindingStorageBufferUpdateAfterBind &&
-            device_info->properties2.properties.limits.minStorageBufferOffsetAlignment <= 4)
+            device_info->properties2.properties.limits.minStorageBufferOffsetAlignment <= 16)
+    {
         flags |= VKD3D_BINDLESS_RAW_SSBO;
 
+        if (device_info->properties2.properties.limits.minStorageBufferOffsetAlignment > 4)
+            flags |= VKD3D_SSBO_OFFSET_BUFFER;
+    }
+
     if (device_info->buffer_device_address_features.bufferDeviceAddress && (flags & VKD3D_BINDLESS_UAV))
         flags |= VKD3D_RAW_VA_UAV_COUNTER;
 
@@ -3469,6 +3494,9 @@ HRESULT vkd3d_bindless_state_init(struct vkd3d_bindless_state *bindless_state,
     if (bindless_state->flags & VKD3D_RAW_VA_UAV_COUNTER)
         extra_bindings |= VKD3D_BINDLESS_SET_EXTRA_UAV_COUNTER_BUFFER;
 
+    if (bindless_state->flags & VKD3D_SSBO_OFFSET_BUFFER)
+        extra_bindings |= VKD3D_BINDLESS_SET_EXTRA_SSBO_OFFSET_BUFFER;
+
     if (FAILED(hr = vkd3d_bindless_state_add_binding(bindless_state, device,
             VKD3D_BINDLESS_SET_SAMPLER, VK_DESCRIPTOR_TYPE_SAMPLER)))
         goto fail;

libs/vkd3d/vkd3d_private.h

@@ -629,6 +629,7 @@ enum vkd3d_descriptor_flag
     VKD3D_DESCRIPTOR_FLAG_DEFINED     = (1 << 0),
     VKD3D_DESCRIPTOR_FLAG_VIEW        = (1 << 1),
     VKD3D_DESCRIPTOR_FLAG_UAV_COUNTER = (1 << 2),
+    VKD3D_DESCRIPTOR_FLAG_SSBO_OFFSET = (1 << 3),
 };
 
 struct vkd3d_descriptor_binding
@@ -725,11 +726,16 @@ static inline struct d3d12_dsv_desc *d3d12_dsv_desc_from_cpu_handle(D3D12_CPU_DE
 void d3d12_dsv_desc_create_dsv(struct d3d12_dsv_desc *dsv_desc, struct d3d12_device *device,
         struct d3d12_resource *resource, const D3D12_DEPTH_STENCIL_VIEW_DESC *desc);
 
-struct d3d12_descriptor_heap_uav_counters
+struct vkd3d_bound_ssbo_range
 {
-    VkDeviceAddress *data;
-    VkDeviceMemory vk_memory;
-    VkBuffer vk_buffer;
+    uint32_t offset;  /* offset to first byte */
+    uint32_t length;  /* bound size in bytes */
+};
+
+struct vkd3d_host_visible_buffer_range
+{
+    VkDescriptorBufferInfo descriptor;
+    void *host_ptr;
 };
 
 /* ID3D12DescriptorHeap */
@@ -743,7 +749,13 @@ struct d3d12_descriptor_heap
     VkDescriptorPool vk_descriptor_pool;
     VkDescriptorSet vk_descriptor_sets[VKD3D_MAX_BINDLESS_DESCRIPTOR_SETS];
 
-    struct d3d12_descriptor_heap_uav_counters uav_counters;
+    VkDeviceMemory vk_memory;
+    VkBuffer vk_buffer;
+    void *host_memory;
+
+    struct vkd3d_host_visible_buffer_range uav_counters;
+    struct vkd3d_host_visible_buffer_range ssbo_ranges;
+
     struct d3d12_device *device;
 
     struct vkd3d_private_store private_store;
@@ -804,6 +816,7 @@ enum vkd3d_root_signature_flag
     VKD3D_ROOT_SIGNATURE_USE_PUSH_DESCRIPTORS       = 0x00000001u,
     VKD3D_ROOT_SIGNATURE_USE_INLINE_UNIFORM_BLOCK   = 0x00000002u,
     VKD3D_ROOT_SIGNATURE_USE_RAW_VA_UAV_COUNTERS    = 0x00000004u,
+    VKD3D_ROOT_SIGNATURE_USE_SSBO_OFFSET_BUFFER     = 0x00000008u,
 };
 
 struct d3d12_root_descriptor_table
@@ -872,6 +885,7 @@ struct d3d12_root_signature
     VkPushConstantRange push_constant_range;
     struct vkd3d_shader_descriptor_binding push_constant_ubo_binding;
     struct vkd3d_shader_descriptor_binding uav_counter_binding;
+    struct vkd3d_shader_descriptor_binding offset_buffer_binding;
 
     uint32_t descriptor_table_offset;
     uint32_t descriptor_table_count;
@@ -1574,6 +1588,7 @@ enum vkd3d_bindless_flags
     VKD3D_RAW_VA_UAV_COUNTER    = (1u << 4),
     VKD3D_BINDLESS_CBV_AS_SSBO  = (1u << 5),
     VKD3D_BINDLESS_RAW_SSBO     = (1u << 6),
+    VKD3D_SSBO_OFFSET_BUFFER    = (1u << 7),
 };
 
 #define VKD3D_BINDLESS_SET_MAX_EXTRA_BINDINGS 8
@@ -1590,6 +1605,7 @@ enum vkd3d_bindless_set_flag
     VKD3D_BINDLESS_SET_RAW_SSBO = (1u << 7),
 
     VKD3D_BINDLESS_SET_EXTRA_UAV_COUNTER_BUFFER = (1u << 24),
+    VKD3D_BINDLESS_SET_EXTRA_SSBO_OFFSET_BUFFER = (1u << 25),
     VKD3D_BINDLESS_SET_EXTRA_MASK = 0xff000000u
 };
 
@@ -1985,6 +2001,14 @@ static inline VkDeviceSize d3d12_device_get_ssbo_alignment(struct d3d12_device *
     return device->device_info.properties2.properties.limits.minStorageBufferOffsetAlignment;
 }
 
+static inline bool d3d12_device_use_ssbo_root_descriptors(struct d3d12_device *device)
+{
+    /* We only know the VA of root SRV/UAVs, so we cannot
+     * make any better assumptions about the alignment */
+    return d3d12_device_use_ssbo_raw_buffer(device) &&
+            d3d12_device_get_ssbo_alignment(device) <= 4;
+}
+
 /* ID3DBlob */
 struct d3d_blob
 {

subprojects/dxil-spirv

@@ -1 +1 @@
-Subproject commit 64ede36e58ebecce64b2f44b908503a8e4ecea19
+Subproject commit 1f85c7bff8cf6b5defc00e300e1d4540d81ac8d1