WIP: Highly experimental evil descriptor hoisting

include/vkd3d.h

@@ -101,6 +101,7 @@ extern "C" {
 #define VKD3D_CONFIG_FLAG_DISABLE_UAV_COMPRESSION (1ull << 43)
 #define VKD3D_CONFIG_FLAG_DISABLE_DEPTH_COMPRESSION (1ull << 44)
 #define VKD3D_CONFIG_FLAG_DISABLE_COLOR_COMPRESSION (1ull << 45)
+#define VKD3D_CONFIG_FLAG_DESCRIPTOR_HOISTING (1ull << 46)
 
 struct vkd3d_instance;
 

include/vkd3d_shader.h

@@ -77,15 +77,28 @@ enum vkd3d_shader_meta_flags
     VKD3D_SHADER_META_FLAG_FORCE_COMPUTE_BARRIER_AFTER_DISPATCH = 1 << 16,
 };
 
+/* Tweaked to make the meta struct a nice and round 64 bytes. */
+#define VKD3D_MAX_HOISTED_DESCRIPTORS 15
+#define VKD3D_MAX_HOISTED_DESCRIPTOR_CONSTANT_OFFSET (1 << 7)
+
+struct vkd3d_shader_meta_hoisted_desc
+{
+    uint16_t table_index : 6; /* API limit */
+    uint16_t constant_offset : 7; /* 128 offset from table should be more than enough ... */
+    uint16_t vk_descriptor_type : 3; /* The standard descriptor types we care about fit into 3 bits as-is. */
+};
+
 struct vkd3d_shader_meta
 {
     vkd3d_shader_hash_t hash;
     unsigned int cs_workgroup_size[3]; /* Only contains valid data if uses_subgroup_size is true. */
     unsigned int patch_vertex_count; /* Relevant for HS. May be 0, in which case the patch vertex count is not known. */
     unsigned int cs_required_wave_size; /* If non-zero, force a specific CS subgroup size. */
     uint32_t flags; /* vkd3d_shader_meta_flags */
+    struct vkd3d_shader_meta_hoisted_desc hoist_desc[VKD3D_MAX_HOISTED_DESCRIPTORS];
+    uint16_t num_hoisted_descriptors;
 };
-STATIC_ASSERT(sizeof(struct vkd3d_shader_meta) == 32);
+STATIC_ASSERT(sizeof(struct vkd3d_shader_meta) == 64);
 
 struct vkd3d_shader_code
 {
@@ -220,6 +233,7 @@ enum vkd3d_shader_interface_flag
     VKD3D_SHADER_INTERFACE_DESCRIPTOR_QA_BUFFER             = 0x00000010u,
     /* In this model, use descriptor_size_cbv_srv_uav as array stride for raw VA buffer. */
     VKD3D_SHADER_INTERFACE_RAW_VA_ALIAS_DESCRIPTOR_BUFFER   = 0x00000020u,
+    VKD3D_SHADER_INTERFACE_HOIST_DESCRIPTORS                = 0x00000040u,
 };
 
 struct vkd3d_shader_stage_io_entry
@@ -278,6 +292,7 @@ struct vkd3d_shader_interface_info
     /* Used for either VKD3D_SHADER_INTERFACE_RAW_VA_ALIAS_DESCRIPTOR_BUFFER or local root signatures. */
     uint32_t descriptor_size_cbv_srv_uav;
     uint32_t descriptor_size_sampler;
+    uint32_t hoist_descriptor_set_index;
 };
 
 struct vkd3d_shader_descriptor_table

libs/vkd3d-shader/dxil.c

@@ -123,6 +123,7 @@ struct vkd3d_dxil_remap_userdata
     const struct vkd3d_shader_interface_info *shader_interface_info;
     const struct vkd3d_shader_interface_local_info *shader_interface_local_info;
     unsigned int num_root_descriptors;
+    struct vkd3d_shader_meta *meta;
 };
 
 struct vkd3d_dxil_remap_info
@@ -253,6 +254,32 @@ static dxil_spv_bool dxil_remap(const struct vkd3d_dxil_remap_userdata *remap,
         return DXIL_SPV_TRUE;
 }
 
+static void dxil_remap_check_hoist(
+        const struct vkd3d_shader_interface_info *shader_interface_info,
+        const struct vkd3d_dxil_remap_userdata *remap, const dxil_spv_d3d_binding *d3d_binding,
+        dxil_spv_vulkan_binding *vk_binding,
+        VkDescriptorType vk_descriptor_type)
+{
+    struct vkd3d_shader_meta_hoisted_desc *hoisted;
+
+    if ((shader_interface_info->flags & VKD3D_SHADER_INTERFACE_HOIST_DESCRIPTORS) &&
+            vk_binding->bindless.use_heap == DXIL_SPV_TRUE &&
+            remap->meta->num_hoisted_descriptors < VKD3D_MAX_HOISTED_DESCRIPTORS &&
+            d3d_binding->range_size == 1 &&
+            vk_binding->bindless.heap_root_offset < VKD3D_MAX_HOISTED_DESCRIPTOR_CONSTANT_OFFSET)
+    {
+        hoisted = &remap->meta->hoist_desc[remap->meta->num_hoisted_descriptors];
+        hoisted->table_index = vk_binding->root_constant_index + remap->num_root_descriptors * 2 -
+                shader_interface_info->descriptor_tables.offset / sizeof(uint32_t);
+        hoisted->constant_offset = vk_binding->bindless.heap_root_offset;
+        hoisted->vk_descriptor_type = vk_descriptor_type;
+        vk_binding->bindless.use_heap = DXIL_SPV_FALSE;
+        vk_binding->set = shader_interface_info->hoist_descriptor_set_index;
+        vk_binding->binding = remap->meta->num_hoisted_descriptors;
+        remap->meta->num_hoisted_descriptors++;
+    }
+}
+
 static dxil_spv_bool dxil_srv_remap(void *userdata, const dxil_spv_d3d_binding *d3d_binding,
                                     dxil_spv_srv_vulkan_binding *vk_binding)
 {
@@ -275,6 +302,12 @@ static dxil_spv_bool dxil_srv_remap(void *userdata, const dxil_spv_d3d_binding *
             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
+        {
+            dxil_remap_check_hoist(shader_interface_info, remap, d3d_binding, &vk_binding->buffer_binding,
+                    VK_DESCRIPTOR_TYPE_STORAGE_BUFFER);
+        }
+
         return DXIL_SPV_TRUE;
     }
     else
@@ -287,16 +320,31 @@ static dxil_spv_bool dxil_srv_remap(void *userdata, const dxil_spv_d3d_binding *
         }
     }
 
-    return dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_SRV,
-            d3d_binding, &vk_binding->buffer_binding, resource_flags);
+    if (!dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_SRV,
+            d3d_binding, &vk_binding->buffer_binding, resource_flags))
+        return DXIL_SPV_FALSE;
+
+    if (!(shader_interface_info->flags & VKD3D_SHADER_INTERFACE_TYPED_OFFSET_BUFFER))
+    {
+        dxil_remap_check_hoist(shader_interface_info, remap, d3d_binding, &vk_binding->buffer_binding,
+                d3d_binding->kind == DXIL_SPV_RESOURCE_KIND_TYPED_BUFFER ?
+                        VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER : VK_DESCRIPTOR_TYPE_SAMPLED_IMAGE);
+    }
+
+    return DXIL_SPV_TRUE;
 }
 
 static dxil_spv_bool dxil_sampler_remap(void *userdata, const dxil_spv_d3d_binding *d3d_binding,
                                         dxil_spv_vulkan_binding *vk_binding)
 {
     const struct vkd3d_dxil_remap_userdata *remap = userdata;
-    return dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_SAMPLER,
-            d3d_binding, vk_binding, VKD3D_SHADER_BINDING_FLAG_IMAGE);
+
+    if (!dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_SAMPLER,
+            d3d_binding, vk_binding, VKD3D_SHADER_BINDING_FLAG_IMAGE))
+        return DXIL_SPV_FALSE;
+
+    dxil_remap_check_hoist(remap->shader_interface_info, remap, d3d_binding, vk_binding, VK_DESCRIPTOR_TYPE_SAMPLER);
+    return DXIL_SPV_TRUE;
 }
 
 static dxil_spv_bool dxil_input_remap(void *userdata, const dxil_spv_d3d_vertex_input *d3d_input,
@@ -421,6 +469,12 @@ static dxil_spv_bool dxil_uav_remap(void *userdata, const dxil_spv_uav_d3d_bindi
                 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 (!d3d_binding->has_counter)
+            {
+                dxil_remap_check_hoist(shader_interface_info, remap,
+                        &d3d_binding->d3d_binding, &vk_binding->buffer_binding,
+                        VK_DESCRIPTOR_TYPE_STORAGE_BUFFER);
+            }
         }
         else if (!dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_UAV, &d3d_binding->d3d_binding,
                 &vk_binding->buffer_binding, resource_flags))
@@ -453,6 +507,13 @@ static dxil_spv_bool dxil_uav_remap(void *userdata, const dxil_spv_uav_d3d_bindi
             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 (!d3d_binding->has_counter)
+        {
+            dxil_remap_check_hoist(shader_interface_info, remap,
+                    &d3d_binding->d3d_binding, &vk_binding->buffer_binding,
+                    d3d_binding->d3d_binding.kind == DXIL_SPV_RESOURCE_KIND_TYPED_BUFFER ?
+                            VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER : VK_DESCRIPTOR_TYPE_STORAGE_IMAGE);
+        }
     }
 
     if (d3d_binding->has_counter)
@@ -499,9 +560,14 @@ static dxil_spv_bool dxil_cbv_remap(void *userdata, const dxil_spv_d3d_binding *
     }
 
     vk_binding->push_constant = DXIL_SPV_FALSE;
-    return dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_CBV,
+    if (!dxil_remap(remap, VKD3D_SHADER_DESCRIPTOR_TYPE_CBV,
             d3d_binding, &vk_binding->vulkan.uniform_binding,
-            VKD3D_SHADER_BINDING_FLAG_BUFFER);
+            VKD3D_SHADER_BINDING_FLAG_BUFFER))
+        return DXIL_SPV_FALSE;
+
+    dxil_remap_check_hoist(remap->shader_interface_info, remap, d3d_binding,
+            &vk_binding->vulkan.uniform_binding, VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER);
+    return DXIL_SPV_TRUE;
 }
 
 static void vkd3d_dxil_log_callback(void *userdata, dxil_spv_log_level level, const char *msg)
@@ -1011,6 +1077,7 @@ int vkd3d_shader_compile_dxil(const struct vkd3d_shader_code *dxbc,
     remap_userdata.shader_interface_info = shader_interface_info;
     remap_userdata.shader_interface_local_info = NULL;
     remap_userdata.num_root_descriptors = num_root_descriptors;
+    remap_userdata.meta = &spirv->meta;
 
     dxil_spv_converter_set_root_constant_word_count(converter, root_constant_words);
     dxil_spv_converter_set_root_descriptor_count(converter, num_root_descriptors);
@@ -1564,6 +1631,7 @@ int vkd3d_shader_compile_dxil_export(const struct vkd3d_shader_code *dxil,
     remap_userdata.shader_interface_info = shader_interface_info;
     remap_userdata.shader_interface_local_info = shader_interface_local_info;
     remap_userdata.num_root_descriptors = num_root_descriptors;
+    remap_userdata.meta = &spirv->meta;
 
     dxil_spv_converter_set_root_constant_word_count(converter, root_constant_words);
     dxil_spv_converter_set_root_descriptor_count(converter, num_root_descriptors);

libs/vkd3d/cache.c

@@ -181,7 +181,7 @@ VkResult vkd3d_create_pipeline_cache(struct d3d12_device *device,
     return VK_CALL(vkCreatePipelineCache(device->vk_device, &info, NULL, cache));
 }
 
-#define VKD3D_CACHE_BLOB_VERSION MAKE_MAGIC('V','K','B',3)
+#define VKD3D_CACHE_BLOB_VERSION MAKE_MAGIC('V','K','B',4)
 
 enum vkd3d_pipeline_blob_chunk_type
 {
@@ -1314,8 +1314,8 @@ struct vkd3d_serialized_pipeline_toc_entry
 };
 STATIC_ASSERT(sizeof(struct vkd3d_serialized_pipeline_toc_entry) == 16);
 
-#define VKD3D_PIPELINE_LIBRARY_VERSION_TOC MAKE_MAGIC('V','K','L',4)
-#define VKD3D_PIPELINE_LIBRARY_VERSION_STREAM MAKE_MAGIC('V','K','S',4)
+#define VKD3D_PIPELINE_LIBRARY_VERSION_TOC MAKE_MAGIC('V','K','L',5)
+#define VKD3D_PIPELINE_LIBRARY_VERSION_STREAM MAKE_MAGIC('V','K','S',5)
 
 struct vkd3d_serialized_pipeline_library_toc
 {