[PATCH vkd3d 05/13] vkd3d: Introduce meta ops for shader-based UAV clears.
Philip Rebohle
philip.rebohle at tu-dortmund.de
Mon Nov 11 10:03:40 CST 2019
Signed-off-by: Philip Rebohle <philip.rebohle at tu-dortmund.de>
---
Makefile.am | 2 +
libs/vkd3d/device.c | 6 +
libs/vkd3d/meta.c | 325 +++++++++++++
libs/vkd3d/vkd3d_private.h | 62 +++
libs/vkd3d/vkd3d_spv_shaders.h | 830 +++++++++++++++++++++++++++++++++
5 files changed, 1225 insertions(+)
create mode 100644 libs/vkd3d/meta.c
create mode 100644 libs/vkd3d/vkd3d_spv_shaders.h
diff --git a/Makefile.am b/Makefile.am
index 2b0e8f3..75d2a97 100644
--- a/Makefile.am
+++ b/Makefile.am
@@ -102,12 +102,14 @@ libvkd3d_la_SOURCES = \
include/vkd3d_unknown.idl \
libs/vkd3d/command.c \
libs/vkd3d/device.c \
+ libs/vkd3d/meta.c \
libs/vkd3d/resource.c \
libs/vkd3d/state.c \
libs/vkd3d/utils.c \
libs/vkd3d/vkd3d.map \
libs/vkd3d/vkd3d_main.c \
libs/vkd3d/vkd3d_private.h \
+ libs/vkd3d/vkd3d_spv_shaders.h \
libs/vkd3d/vulkan_procs.h \
libs/vkd3d_version.c
libvkd3d_la_LDFLAGS = $(AM_LDFLAGS) -version-info 2:0:1
diff --git a/libs/vkd3d/device.c b/libs/vkd3d/device.c
index 0624318..114d671 100644
--- a/libs/vkd3d/device.c
+++ b/libs/vkd3d/device.c
@@ -2154,6 +2154,7 @@ static ULONG STDMETHODCALLTYPE d3d12_device_Release(ID3D12Device *iface)
vkd3d_private_store_destroy(&device->private_store);
vkd3d_cleanup_format_info(device);
+ vkd3d_meta_ops_destroy(&device->meta_ops, device);
vkd3d_destroy_null_resources(&device->null_resources, device);
vkd3d_gpu_va_allocator_cleanup(&device->gpu_va_allocator);
vkd3d_render_pass_cache_cleanup(&device->render_pass_cache, device);
@@ -3447,6 +3448,9 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
if (FAILED(hr = vkd3d_init_null_resources(&device->null_resources, device)))
goto out_cleanup_format_info;
+ if (FAILED(hr = vkd3d_meta_ops_init(&device->meta_ops, device)))
+ goto out_destroy_null_resources;
+
vkd3d_render_pass_cache_init(&device->render_pass_cache);
vkd3d_gpu_va_allocator_init(&device->gpu_va_allocator);
@@ -3458,6 +3462,8 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
return S_OK;
+out_destroy_null_resources:
+ vkd3d_destroy_null_resources(&device->null_resources, device);
out_cleanup_format_info:
vkd3d_cleanup_format_info(device);
out_stop_fence_worker:
diff --git a/libs/vkd3d/meta.c b/libs/vkd3d/meta.c
new file mode 100644
index 0000000..9a18cf6
--- /dev/null
+++ b/libs/vkd3d/meta.c
@@ -0,0 +1,325 @@
+#include "vkd3d_private.h"
+#include "vkd3d_spv_shaders.h"
+
+#define SPIRV_CODE(name) name, sizeof(name)
+
+static VkResult vkd3d_create_shader_module(struct d3d12_device *device,
+ size_t code_size, const uint32_t *code, VkShaderModule *module)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ VkShaderModuleCreateInfo shader_module_info;
+
+ shader_module_info.sType = VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO;
+ shader_module_info.pNext = NULL;
+ shader_module_info.flags = 0;
+ shader_module_info.codeSize = code_size;
+ shader_module_info.pCode = code;
+
+ return VK_CALL(vkCreateShaderModule(device->vk_device, &shader_module_info, NULL, module));
+}
+
+static VkResult vkd3d_create_descriptor_set_layout(struct d3d12_device *device,
+ uint32_t binding_count, const VkDescriptorSetLayoutBinding *bindings, VkDescriptorSetLayout *set_layout)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ VkDescriptorSetLayoutCreateInfo set_layout_info;
+
+ set_layout_info.sType = VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO;
+ set_layout_info.pNext = NULL;
+ set_layout_info.flags = 0;
+ set_layout_info.bindingCount = binding_count;
+ set_layout_info.pBindings = bindings;
+
+ return VK_CALL(vkCreateDescriptorSetLayout(device->vk_device, &set_layout_info, NULL, set_layout));
+}
+
+static VkResult vkd3d_create_pipeline_layout(struct d3d12_device *device,
+ uint32_t set_layout_count, const VkDescriptorSetLayout *set_layouts,
+ uint32_t push_constant_range_count, const VkPushConstantRange *push_constant_ranges,
+ VkPipelineLayout *pipeline_layout)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ VkPipelineLayoutCreateInfo pipeline_layout_info;
+
+ pipeline_layout_info.sType = VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO;
+ pipeline_layout_info.pNext = NULL;
+ pipeline_layout_info.flags = 0;
+ pipeline_layout_info.setLayoutCount = set_layout_count;
+ pipeline_layout_info.pSetLayouts = set_layouts;
+ pipeline_layout_info.pushConstantRangeCount = push_constant_range_count;
+ pipeline_layout_info.pPushConstantRanges = push_constant_ranges;
+
+ return VK_CALL(vkCreatePipelineLayout(device->vk_device, &pipeline_layout_info, NULL, pipeline_layout));
+}
+
+static VkResult vkd3d_create_compute_pipeline(struct d3d12_device *device,
+ size_t code_size, const uint32_t *code, VkPipelineLayout layout,
+ const VkSpecializationInfo *specialization_info, VkPipeline *pipeline)
+{
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ VkComputePipelineCreateInfo pipeline_info;
+ VkResult vr;
+
+ pipeline_info.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO;
+ pipeline_info.pNext = NULL;
+ pipeline_info.flags = 0;
+ pipeline_info.layout = layout;
+ pipeline_info.basePipelineHandle = VK_NULL_HANDLE;
+ pipeline_info.basePipelineIndex = -1;
+
+ pipeline_info.stage.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO;
+ pipeline_info.stage.pNext = NULL;
+ pipeline_info.stage.flags = 0;
+ pipeline_info.stage.stage = VK_SHADER_STAGE_COMPUTE_BIT;
+ pipeline_info.stage.pName = "main";
+ pipeline_info.stage.pSpecializationInfo = specialization_info;
+
+ if ((vr = vkd3d_create_shader_module(device, code_size, code, &pipeline_info.stage.module)) < 0)
+ {
+ ERR("Failed to create shader module, vr %d.", vr);
+ return vr;
+ }
+
+ vr = VK_CALL(vkCreateComputePipelines(device->vk_device, VK_NULL_HANDLE, 1, &pipeline_info, NULL, pipeline));
+ VK_CALL(vkDestroyShaderModule(device->vk_device, pipeline_info.stage.module, NULL));
+
+ return vr;
+}
+
+HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ struct d3d12_device *device)
+{
+ VkDescriptorSetLayoutBinding set_binding;
+ VkPushConstantRange push_constant_range;
+ unsigned int i;
+ VkResult vr;
+
+ struct {
+ VkDescriptorSetLayout *set_layout;
+ VkPipelineLayout *pipeline_layout;
+ VkDescriptorType descriptor_type;
+ }
+ set_layouts[] =
+ {
+ { &meta_clear_uav_ops->vk_set_layout_buffer, &meta_clear_uav_ops->vk_pipeline_layout_buffer, VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER },
+ { &meta_clear_uav_ops->vk_set_layout_image, &meta_clear_uav_ops->vk_pipeline_layout_image, VK_DESCRIPTOR_TYPE_STORAGE_IMAGE },
+ };
+
+ struct {
+ VkPipeline *pipeline;
+ VkPipelineLayout *pipeline_layout;
+ const uint32_t *code;
+ size_t code_size;
+ }
+ pipelines[] =
+ {
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_buffer,
+ &meta_clear_uav_ops->vk_pipeline_layout_buffer,
+ SPIRV_CODE(cs_clear_uav_buffer_float_spv) },
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_1d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_1d_float_spv) },
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_1d_array,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_1d_array_float_spv) },
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_2d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_2d_float_spv) },
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_2d_array,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_2d_array_float_spv) },
+ { &meta_clear_uav_ops->clear_float.vk_pipeline_clear_image_3d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_3d_float_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_buffer,
+ &meta_clear_uav_ops->vk_pipeline_layout_buffer,
+ SPIRV_CODE(cs_clear_uav_buffer_uint_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_1d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_1d_uint_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_1d_array,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_1d_array_uint_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_2d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_2d_uint_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_2d_array,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_2d_array_uint_spv) },
+ { &meta_clear_uav_ops->clear_uint.vk_pipeline_clear_image_3d,
+ &meta_clear_uav_ops->vk_pipeline_layout_image,
+ SPIRV_CODE(cs_clear_uav_image_3d_uint_spv) },
+ };
+
+ memset(meta_clear_uav_ops, 0, sizeof(*meta_clear_uav_ops));
+
+ set_binding.binding = 0;
+ set_binding.descriptorCount = 1;
+ set_binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+ set_binding.pImmutableSamplers = NULL;
+
+ push_constant_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
+ push_constant_range.offset = 0;
+ push_constant_range.size = sizeof(struct vkd3d_clear_uav_args);
+
+ for (i = 0; i < ARRAY_SIZE(set_layouts); i++)
+ {
+ set_binding.descriptorType = set_layouts[i].descriptor_type;
+
+ vr = vkd3d_create_descriptor_set_layout(device, 1, &set_binding, set_layouts[i].set_layout);
+
+ if (vr < 0)
+ {
+ ERR("Failed to create descriptor set layout %u, vr %d.", i, vr);
+ goto fail;
+ }
+
+ vr = vkd3d_create_pipeline_layout(device, 1, set_layouts[i].set_layout,
+ 1, &push_constant_range, set_layouts[i].pipeline_layout);
+
+ if (vr < 0)
+ {
+ ERR("Failed to create pipeline layout %u, vr %d.", i, vr);
+ goto fail;
+ }
+ }
+
+ for (i = 0; i < ARRAY_SIZE(pipelines); i++)
+ {
+ if ((vr = vkd3d_create_compute_pipeline(device, pipelines[i].code_size, pipelines[i].code,
+ *pipelines[i].pipeline_layout, NULL, pipelines[i].pipeline)) < 0)
+ {
+ ERR("Failed to create compute pipeline %u, vr %d.", i, vr);
+ goto fail;
+ }
+ }
+
+ return S_OK;
+fail:
+ vkd3d_clear_uav_ops_destroy(meta_clear_uav_ops, device);
+ return hresult_from_vk_result(vr);
+}
+
+void vkd3d_clear_uav_ops_destroy(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ struct d3d12_device *device) {
+ const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
+ unsigned int i;
+
+ struct vkd3d_clear_uav_pipelines* pipeline_sets[] =
+ {
+ &meta_clear_uav_ops->clear_float,
+ &meta_clear_uav_ops->clear_uint,
+ };
+
+ VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, meta_clear_uav_ops->vk_set_layout_buffer, NULL));
+ VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, meta_clear_uav_ops->vk_set_layout_image, NULL));
+
+ VK_CALL(vkDestroyPipelineLayout(device->vk_device, meta_clear_uav_ops->vk_pipeline_layout_buffer, NULL));
+ VK_CALL(vkDestroyPipelineLayout(device->vk_device, meta_clear_uav_ops->vk_pipeline_layout_image, NULL));
+
+ for (i = 0; i < ARRAY_SIZE(pipeline_sets); i++)
+ {
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_buffer, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_1d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_2d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_3d, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_1d_array, NULL));
+ VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->vk_pipeline_clear_image_2d_array, NULL));
+ }
+}
+
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_buffer_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ bool as_uint)
+{
+ struct vkd3d_clear_uav_pipeline info;
+
+ const struct vkd3d_clear_uav_pipelines *pipelines = as_uint
+ ? &meta_clear_uav_ops->clear_uint
+ : &meta_clear_uav_ops->clear_float;
+
+ info.vk_set_layout = meta_clear_uav_ops->vk_set_layout_buffer;
+ info.vk_pipeline_layout = meta_clear_uav_ops->vk_pipeline_layout_buffer;
+ info.vk_pipeline = pipelines->vk_pipeline_clear_buffer;
+ return info;
+}
+
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_image_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ VkImageViewType image_view_type, bool as_uint)
+{
+ struct vkd3d_clear_uav_pipeline info;
+
+ const struct vkd3d_clear_uav_pipelines *pipelines = as_uint
+ ? &meta_clear_uav_ops->clear_uint
+ : &meta_clear_uav_ops->clear_float;
+
+ info.vk_set_layout = meta_clear_uav_ops->vk_set_layout_image;
+ info.vk_pipeline_layout = meta_clear_uav_ops->vk_pipeline_layout_image;
+
+ switch (image_view_type)
+ {
+ case VK_IMAGE_VIEW_TYPE_1D:
+ info.vk_pipeline = pipelines->vk_pipeline_clear_image_1d;
+ break;
+ case VK_IMAGE_VIEW_TYPE_2D:
+ info.vk_pipeline = pipelines->vk_pipeline_clear_image_2d;
+ break;
+ case VK_IMAGE_VIEW_TYPE_3D:
+ info.vk_pipeline = pipelines->vk_pipeline_clear_image_3d;
+ break;
+ case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+ info.vk_pipeline = pipelines->vk_pipeline_clear_image_1d_array;
+ break;
+ case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+ info.vk_pipeline = pipelines->vk_pipeline_clear_image_2d_array;
+ break;
+ default:
+ ERR("Unhandled view type %d.\n", image_view_type);
+ info.vk_pipeline = VK_NULL_HANDLE;
+ }
+
+ return info;
+}
+
+VkExtent3D vkd3d_get_clear_image_uav_workgroup_size(VkImageViewType view_type)
+{
+ switch (view_type)
+ {
+ case VK_IMAGE_VIEW_TYPE_1D:
+ case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
+ {
+ VkExtent3D result = { 64, 1, 1 };
+ return result;
+ }
+ case VK_IMAGE_VIEW_TYPE_2D:
+ case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
+ case VK_IMAGE_VIEW_TYPE_3D:
+ {
+ VkExtent3D result = { 8, 8, 1 };
+ return result;
+ }
+ default:
+ {
+ VkExtent3D result = { 0, 0, 0 };
+ ERR("Unhandled view type %d.\n", view_type);
+ return result;
+ }
+ }
+}
+
+HRESULT vkd3d_meta_ops_init(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device)
+{
+ HRESULT hr;
+
+ memset(meta_ops, 0, sizeof(*meta_ops));
+
+ if (FAILED(hr = vkd3d_clear_uav_ops_init(&meta_ops->clear_uav, device)))
+ return hr;
+
+ return S_OK;
+}
+
+HRESULT vkd3d_meta_ops_destroy(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device)
+{
+ vkd3d_clear_uav_ops_destroy(&meta_ops->clear_uav, device);
+ return S_OK;
+}
diff --git a/libs/vkd3d/vkd3d_private.h b/libs/vkd3d/vkd3d_private.h
index 84b5ff2..d0224d3 100644
--- a/libs/vkd3d/vkd3d_private.h
+++ b/libs/vkd3d/vkd3d_private.h
@@ -1059,6 +1059,67 @@ struct vkd3d_format_compatibility_list
VkFormat vk_formats[VKD3D_MAX_COMPATIBLE_FORMAT_COUNT];
};
+/* meta operations */
+struct vkd3d_clear_uav_args
+{
+ VkClearColorValue clear_color;
+ VkOffset2D offset;
+ VkExtent2D extent;
+};
+
+struct vkd3d_clear_uav_pipelines
+{
+ VkPipeline vk_pipeline_clear_buffer;
+ VkPipeline vk_pipeline_clear_image_1d;
+ VkPipeline vk_pipeline_clear_image_2d;
+ VkPipeline vk_pipeline_clear_image_3d;
+ VkPipeline vk_pipeline_clear_image_1d_array;
+ VkPipeline vk_pipeline_clear_image_2d_array;
+};
+
+struct vkd3d_clear_uav_ops
+{
+ VkDescriptorSetLayout vk_set_layout_buffer;
+ VkDescriptorSetLayout vk_set_layout_image;
+
+ VkPipelineLayout vk_pipeline_layout_buffer;
+ VkPipelineLayout vk_pipeline_layout_image;
+
+ struct vkd3d_clear_uav_pipelines clear_float;
+ struct vkd3d_clear_uav_pipelines clear_uint;
+};
+
+struct vkd3d_clear_uav_pipeline
+{
+ VkDescriptorSetLayout vk_set_layout;
+ VkPipelineLayout vk_pipeline_layout;
+ VkPipeline vk_pipeline;
+};
+
+HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ struct d3d12_device *device) DECLSPEC_HIDDEN;
+void vkd3d_clear_uav_ops_destroy(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ struct d3d12_device *device) DECLSPEC_HIDDEN;
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_buffer_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ bool as_uint) DECLSPEC_HIDDEN;
+struct vkd3d_clear_uav_pipeline vkd3d_clear_uav_ops_get_clear_image_pipeline(const struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
+ VkImageViewType image_view_type, bool as_uint) DECLSPEC_HIDDEN;
+VkExtent3D vkd3d_get_clear_image_uav_workgroup_size(VkImageViewType view_type) DECLSPEC_HIDDEN;
+
+inline VkExtent3D vkd3d_get_clear_buffer_uav_workgroup_size()
+{
+ VkExtent3D result = { 128, 1, 1 };
+ return result;
+}
+
+struct vkd3d_meta_ops
+{
+ struct vkd3d_clear_uav_ops clear_uav;
+};
+
+HRESULT vkd3d_meta_ops_init(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device) DECLSPEC_HIDDEN;
+HRESULT vkd3d_meta_ops_destroy(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device) DECLSPEC_HIDDEN;
+
/* ID3D12Device */
struct d3d12_device
{
@@ -1104,6 +1165,7 @@ struct d3d12_device
unsigned int format_compatibility_list_count;
const struct vkd3d_format_compatibility_list *format_compatibility_lists;
struct vkd3d_null_resources null_resources;
+ struct vkd3d_meta_ops meta_ops;
};
HRESULT d3d12_device_create(struct vkd3d_instance *instance,
diff --git a/libs/vkd3d/vkd3d_spv_shaders.h b/libs/vkd3d/vkd3d_spv_shaders.h
new file mode 100644
index 0000000..bf86f6b
--- /dev/null
+++ b/libs/vkd3d/vkd3d_spv_shaders.h
@@ -0,0 +1,830 @@
+#ifndef __VKD3D_SPV_SHADERS_H
+#define __VKD3D_SPV_SHADERS_H
+
+const uint32_t cs_clear_uav_buffer_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 128) in;
+
+layout(binding = 0)
+writeonly uniform imageBuffer dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ int thread_id = int(gl_GlobalInvocationID.x);
+
+ if (thread_id < u_info.dst_extent.x)
+ imageStore(dst, u_info.dst_offset.x + thread_id, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000031,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002f,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000c,0x00060010,0x00000004,0x00000011,0x00000080,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000008,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000c,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000023,0x00747364,0x00040047,0x0000000c,0x0000000b,0x0000001c,0x00050048,0x00000016,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+ 0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+ 0x00030047,0x00000023,0x00000019,0x00040047,0x00000030,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040020,0x00000007,0x00000007,0x00000006,0x00040015,0x00000009,0x00000020,0x00000000,
+ 0x00040017,0x0000000a,0x00000009,0x00000003,0x00040020,0x0000000b,0x00000001,0x0000000a,
+ 0x0004003b,0x0000000b,0x0000000c,0x00000001,0x0004002b,0x00000009,0x0000000d,0x00000000,
+ 0x00040020,0x0000000e,0x00000001,0x00000009,0x00030016,0x00000013,0x00000020,0x00040017,
+ 0x00000014,0x00000013,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,0x0005001e,
+ 0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,0x00000016,
+ 0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,0x00000002,
+ 0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,0x00000021,
+ 0x00000013,0x00000005,0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,
+ 0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,0x0004002b,
+ 0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,0x0000002a,0x00000000,0x00040020,
+ 0x0000002b,0x00000009,0x00000014,0x0004002b,0x00000009,0x0000002e,0x00000080,0x0004002b,
+ 0x00000009,0x0000002f,0x00000001,0x0006002c,0x0000000a,0x00000030,0x0000002e,0x0000002f,
+ 0x0000002f,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,
+ 0x0004003b,0x00000007,0x00000008,0x00000007,0x00050041,0x0000000e,0x0000000f,0x0000000c,
+ 0x0000000d,0x0004003d,0x00000009,0x00000010,0x0000000f,0x0004007c,0x00000006,0x00000011,
+ 0x00000010,0x0003003e,0x00000008,0x00000011,0x0004003d,0x00000006,0x00000012,0x00000008,
+ 0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x0000000d,0x0004003d,0x00000006,
+ 0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001c,0x000300f7,
+ 0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+ 0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,0x00000018,
+ 0x00000025,0x0000000d,0x0004003d,0x00000006,0x00000027,0x00000026,0x0004003d,0x00000006,
+ 0x00000028,0x00000008,0x00050080,0x00000006,0x00000029,0x00000027,0x00000028,0x00050041,
+ 0x0000002b,0x0000002c,0x00000018,0x0000002a,0x0004003d,0x00000014,0x0000002d,0x0000002c,
+ 0x00040063,0x00000024,0x00000029,0x0000002d,0x000200f9,0x00000020,0x000200f8,0x00000020,
+ 0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_buffer_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 128) in;
+
+layout(binding = 0)
+writeonly uniform uimageBuffer dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ int thread_id = int(gl_GlobalInvocationID.x);
+
+ if (thread_id < u_info.dst_extent.x)
+ imageStore(dst, u_info.dst_offset.x + thread_id, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000030,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002f,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000c,0x00060010,0x00000004,0x00000011,0x00000080,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000008,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000c,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000015,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000022,0x00747364,0x00040047,0x0000000c,0x0000000b,0x0000001c,0x00050048,0x00000015,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000015,0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,
+ 0x00040047,0x00000022,0x00000022,0x00000000,0x00040047,0x00000022,0x00000021,0x00000000,
+ 0x00030047,0x00000022,0x00000019,0x00040047,0x0000002f,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040020,0x00000007,0x00000007,0x00000006,0x00040015,0x00000009,0x00000020,0x00000000,
+ 0x00040017,0x0000000a,0x00000009,0x00000003,0x00040020,0x0000000b,0x00000001,0x0000000a,
+ 0x0004003b,0x0000000b,0x0000000c,0x00000001,0x0004002b,0x00000009,0x0000000d,0x00000000,
+ 0x00040020,0x0000000e,0x00000001,0x00000009,0x00040017,0x00000013,0x00000009,0x00000004,
+ 0x00040017,0x00000014,0x00000006,0x00000002,0x0005001e,0x00000015,0x00000013,0x00000014,
+ 0x00000014,0x00040020,0x00000016,0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,
+ 0x00000009,0x0004002b,0x00000006,0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,
+ 0x00000006,0x00020014,0x0000001c,0x00090019,0x00000020,0x00000009,0x00000005,0x00000000,
+ 0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000021,0x00000000,0x00000020,
+ 0x0004003b,0x00000021,0x00000022,0x00000000,0x0004002b,0x00000006,0x00000024,0x00000001,
+ 0x0004002b,0x00000006,0x00000029,0x00000000,0x00040020,0x0000002a,0x00000009,0x00000013,
+ 0x0004002b,0x00000009,0x0000002d,0x00000080,0x0004002b,0x00000009,0x0000002e,0x00000001,
+ 0x0006002c,0x0000000a,0x0000002f,0x0000002d,0x0000002e,0x0000002e,0x00050036,0x00000002,
+ 0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000007,0x00000008,
+ 0x00000007,0x00050041,0x0000000e,0x0000000f,0x0000000c,0x0000000d,0x0004003d,0x00000009,
+ 0x00000010,0x0000000f,0x0004007c,0x00000006,0x00000011,0x00000010,0x0003003e,0x00000008,
+ 0x00000011,0x0004003d,0x00000006,0x00000012,0x00000008,0x00060041,0x00000019,0x0000001a,
+ 0x00000017,0x00000018,0x0000000d,0x0004003d,0x00000006,0x0000001b,0x0000001a,0x000500b1,
+ 0x0000001c,0x0000001d,0x00000012,0x0000001b,0x000300f7,0x0000001f,0x00000000,0x000400fa,
+ 0x0000001d,0x0000001e,0x0000001f,0x000200f8,0x0000001e,0x0004003d,0x00000020,0x00000023,
+ 0x00000022,0x00060041,0x00000019,0x00000025,0x00000017,0x00000024,0x0000000d,0x0004003d,
+ 0x00000006,0x00000026,0x00000025,0x0004003d,0x00000006,0x00000027,0x00000008,0x00050080,
+ 0x00000006,0x00000028,0x00000026,0x00000027,0x00050041,0x0000002a,0x0000002b,0x00000017,
+ 0x00000029,0x0004003d,0x00000013,0x0000002c,0x0000002b,0x00040063,0x00000023,0x00000028,
+ 0x0000002c,0x000200f9,0x0000001f,0x000200f8,0x0000001f,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_array_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform image1DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (thread_id.x < u_info.dst_extent.x)
+ imageStore(dst, ivec2(u_info.dst_offset.x + thread_id.x, thread_id.y), u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000036,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000017,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000017,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000017,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000017,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000019,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000024,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000017,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000017,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000017,0x00000002,0x00000023,0x00000018,0x00030047,0x00000017,0x00000002,
+ 0x00040047,0x00000024,0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,
+ 0x00030047,0x00000024,0x00000019,0x00040047,0x00000035,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+ 0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+ 0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+ 0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+ 0x00030016,0x00000014,0x00000020,0x00040017,0x00000015,0x00000014,0x00000004,0x00040017,
+ 0x00000016,0x00000006,0x00000002,0x0005001e,0x00000017,0x00000015,0x00000016,0x00000016,
+ 0x00040020,0x00000018,0x00000009,0x00000017,0x0004003b,0x00000018,0x00000019,0x00000009,
+ 0x0004002b,0x00000006,0x0000001a,0x00000002,0x00040020,0x0000001b,0x00000009,0x00000006,
+ 0x00020014,0x0000001e,0x00090019,0x00000022,0x00000014,0x00000000,0x00000000,0x00000001,
+ 0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,0x00000022,0x0004003b,
+ 0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,0x00000001,0x0004002b,
+ 0x0000000a,0x0000002c,0x00000001,0x0004002b,0x00000006,0x00000030,0x00000000,0x00040020,
+ 0x00000031,0x00000009,0x00000015,0x0004002b,0x0000000a,0x00000034,0x00000040,0x0006002c,
+ 0x0000000b,0x00000035,0x00000034,0x0000002c,0x0000002c,0x00050036,0x00000002,0x00000004,
+ 0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,
+ 0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,
+ 0x0003003e,0x00000009,0x0000000f,0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,
+ 0x0004003d,0x00000006,0x00000013,0x00000012,0x00060041,0x0000001b,0x0000001c,0x00000019,
+ 0x0000001a,0x00000010,0x0004003d,0x00000006,0x0000001d,0x0000001c,0x000500b1,0x0000001e,
+ 0x0000001f,0x00000013,0x0000001d,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,
+ 0x00000020,0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,
+ 0x00060041,0x0000001b,0x00000027,0x00000019,0x00000026,0x00000010,0x0004003d,0x00000006,
+ 0x00000028,0x00000027,0x00050041,0x00000011,0x00000029,0x00000009,0x00000010,0x0004003d,
+ 0x00000006,0x0000002a,0x00000029,0x00050080,0x00000006,0x0000002b,0x00000028,0x0000002a,
+ 0x00050041,0x00000011,0x0000002d,0x00000009,0x0000002c,0x0004003d,0x00000006,0x0000002e,
+ 0x0000002d,0x00050050,0x00000016,0x0000002f,0x0000002b,0x0000002e,0x00050041,0x00000031,
+ 0x00000032,0x00000019,0x00000030,0x0004003d,0x00000015,0x00000033,0x00000032,0x00040063,
+ 0x00000025,0x0000002f,0x00000033,0x000200f9,0x00000021,0x000200f8,0x00000021,0x000100fd,
+ 0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_array_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform uimage1DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (thread_id.x < u_info.dst_extent.x)
+ imageStore(dst, ivec2(u_info.dst_offset.x + thread_id.x, thread_id.y), u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000035,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000023,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000016,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+ 0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+ 0x00030047,0x00000023,0x00000019,0x00040047,0x00000034,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+ 0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+ 0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+ 0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+ 0x00040017,0x00000014,0x0000000a,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,
+ 0x0005001e,0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,
+ 0x00000016,0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,
+ 0x00000002,0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,
+ 0x00000021,0x0000000a,0x00000000,0x00000000,0x00000001,0x00000000,0x00000002,0x00000000,
+ 0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,
+ 0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x0000000a,0x0000002b,0x00000001,
+ 0x0004002b,0x00000006,0x0000002f,0x00000000,0x00040020,0x00000030,0x00000009,0x00000014,
+ 0x0004002b,0x0000000a,0x00000033,0x00000040,0x0006002c,0x0000000b,0x00000034,0x00000033,
+ 0x0000002b,0x0000002b,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,
+ 0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,
+ 0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,
+ 0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000013,
+ 0x00000012,0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x00000010,0x0004003d,
+ 0x00000006,0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000013,0x0000001c,
+ 0x000300f7,0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,
+ 0x0000001f,0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,
+ 0x00000018,0x00000025,0x00000010,0x0004003d,0x00000006,0x00000027,0x00000026,0x00050041,
+ 0x00000011,0x00000028,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000029,0x00000028,
+ 0x00050080,0x00000006,0x0000002a,0x00000027,0x00000029,0x00050041,0x00000011,0x0000002c,
+ 0x00000009,0x0000002b,0x0004003d,0x00000006,0x0000002d,0x0000002c,0x00050050,0x00000015,
+ 0x0000002e,0x0000002a,0x0000002d,0x00050041,0x00000030,0x00000031,0x00000018,0x0000002f,
+ 0x0004003d,0x00000014,0x00000032,0x00000031,0x00040063,0x00000024,0x0000002e,0x00000032,
+ 0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform image1D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (thread_id.x < u_info.dst_extent.x)
+ imageStore(dst, u_info.dst_offset.x + thread_id.x, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000033,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000017,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000017,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000017,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000017,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000019,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000024,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000017,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000017,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000017,0x00000002,0x00000023,0x00000018,0x00030047,0x00000017,0x00000002,
+ 0x00040047,0x00000024,0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,
+ 0x00030047,0x00000024,0x00000019,0x00040047,0x00000032,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+ 0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+ 0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+ 0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+ 0x00030016,0x00000014,0x00000020,0x00040017,0x00000015,0x00000014,0x00000004,0x00040017,
+ 0x00000016,0x00000006,0x00000002,0x0005001e,0x00000017,0x00000015,0x00000016,0x00000016,
+ 0x00040020,0x00000018,0x00000009,0x00000017,0x0004003b,0x00000018,0x00000019,0x00000009,
+ 0x0004002b,0x00000006,0x0000001a,0x00000002,0x00040020,0x0000001b,0x00000009,0x00000006,
+ 0x00020014,0x0000001e,0x00090019,0x00000022,0x00000014,0x00000000,0x00000000,0x00000000,
+ 0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,0x00000022,0x0004003b,
+ 0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,0x00000001,0x0004002b,
+ 0x00000006,0x0000002c,0x00000000,0x00040020,0x0000002d,0x00000009,0x00000015,0x0004002b,
+ 0x0000000a,0x00000030,0x00000040,0x0004002b,0x0000000a,0x00000031,0x00000001,0x0006002c,
+ 0x0000000b,0x00000032,0x00000030,0x00000031,0x00000031,0x00050036,0x00000002,0x00000004,
+ 0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,
+ 0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,
+ 0x0003003e,0x00000009,0x0000000f,0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,
+ 0x0004003d,0x00000006,0x00000013,0x00000012,0x00060041,0x0000001b,0x0000001c,0x00000019,
+ 0x0000001a,0x00000010,0x0004003d,0x00000006,0x0000001d,0x0000001c,0x000500b1,0x0000001e,
+ 0x0000001f,0x00000013,0x0000001d,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,
+ 0x00000020,0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,
+ 0x00060041,0x0000001b,0x00000027,0x00000019,0x00000026,0x00000010,0x0004003d,0x00000006,
+ 0x00000028,0x00000027,0x00050041,0x00000011,0x00000029,0x00000009,0x00000010,0x0004003d,
+ 0x00000006,0x0000002a,0x00000029,0x00050080,0x00000006,0x0000002b,0x00000028,0x0000002a,
+ 0x00050041,0x0000002d,0x0000002e,0x00000019,0x0000002c,0x0004003d,0x00000015,0x0000002f,
+ 0x0000002e,0x00040063,0x00000025,0x0000002b,0x0000002f,0x000200f9,0x00000021,0x000200f8,
+ 0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_1d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 64) in;
+
+layout(binding = 0)
+writeonly uniform uimage1D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (thread_id.x < u_info.dst_extent.x)
+ imageStore(dst, u_info.dst_offset.x + thread_id.x, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000032,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x0000002c,0x00020011,0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,
+ 0x00000000,0x0003000e,0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,
+ 0x00000000,0x0000000d,0x00060010,0x00000004,0x00000011,0x00000040,0x00000001,0x00000001,
+ 0x00030003,0x00000002,0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,
+ 0x00000009,0x65726874,0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,
+ 0x766e496c,0x7461636f,0x496e6f69,0x00000044,0x00050005,0x00000016,0x6e695f75,0x745f6f66,
+ 0x00000000,0x00060006,0x00000016,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,
+ 0x00000016,0x00000001,0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000016,0x00000002,
+ 0x5f747364,0x65747865,0x0000746e,0x00040005,0x00000018,0x6e695f75,0x00006f66,0x00030005,
+ 0x00000023,0x00747364,0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000016,
+ 0x00000000,0x00000023,0x00000000,0x00050048,0x00000016,0x00000001,0x00000023,0x00000010,
+ 0x00050048,0x00000016,0x00000002,0x00000023,0x00000018,0x00030047,0x00000016,0x00000002,
+ 0x00040047,0x00000023,0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,
+ 0x00030047,0x00000023,0x00000019,0x00040047,0x00000031,0x0000000b,0x00000019,0x00020013,
+ 0x00000002,0x00030021,0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,
+ 0x00040017,0x00000007,0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,
+ 0x00040015,0x0000000a,0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,
+ 0x00040020,0x0000000c,0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,
+ 0x0004002b,0x0000000a,0x00000010,0x00000000,0x00040020,0x00000011,0x00000007,0x00000006,
+ 0x00040017,0x00000014,0x0000000a,0x00000004,0x00040017,0x00000015,0x00000006,0x00000002,
+ 0x0005001e,0x00000016,0x00000014,0x00000015,0x00000015,0x00040020,0x00000017,0x00000009,
+ 0x00000016,0x0004003b,0x00000017,0x00000018,0x00000009,0x0004002b,0x00000006,0x00000019,
+ 0x00000002,0x00040020,0x0000001a,0x00000009,0x00000006,0x00020014,0x0000001d,0x00090019,
+ 0x00000021,0x0000000a,0x00000000,0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,
+ 0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,0x00000023,0x00000000,
+ 0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,0x0000002b,0x00000000,
+ 0x00040020,0x0000002c,0x00000009,0x00000014,0x0004002b,0x0000000a,0x0000002f,0x00000040,
+ 0x0004002b,0x0000000a,0x00000030,0x00000001,0x0006002c,0x0000000b,0x00000031,0x0000002f,
+ 0x00000030,0x00000030,0x00050036,0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,
+ 0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,
+ 0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,
+ 0x00050041,0x00000011,0x00000012,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000013,
+ 0x00000012,0x00060041,0x0000001a,0x0000001b,0x00000018,0x00000019,0x00000010,0x0004003d,
+ 0x00000006,0x0000001c,0x0000001b,0x000500b1,0x0000001d,0x0000001e,0x00000013,0x0000001c,
+ 0x000300f7,0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,
+ 0x0000001f,0x0004003d,0x00000021,0x00000024,0x00000023,0x00060041,0x0000001a,0x00000026,
+ 0x00000018,0x00000025,0x00000010,0x0004003d,0x00000006,0x00000027,0x00000026,0x00050041,
+ 0x00000011,0x00000028,0x00000009,0x00000010,0x0004003d,0x00000006,0x00000029,0x00000028,
+ 0x00050080,0x00000006,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+ 0x00000018,0x0000002b,0x0004003d,0x00000014,0x0000002e,0x0000002d,0x00040063,0x00000024,
+ 0x0000002a,0x0000002e,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_array_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform image2DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+ imageStore(dst, ivec3(u_info.dst_offset.xy + thread_id.xy, thread_id.z), u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x0000003a,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+ 0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+ 0x00000019,0x00040047,0x00000039,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+ 0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+ 0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+ 0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+ 0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000001,
+ 0x00000000,0x00000001,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+ 0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+ 0x00000001,0x0004002b,0x0000000a,0x0000002c,0x00000002,0x00040020,0x0000002d,0x00000007,
+ 0x00000006,0x0004002b,0x00000006,0x00000033,0x00000000,0x00040020,0x00000034,0x00000009,
+ 0x00000014,0x0004002b,0x0000000a,0x00000037,0x00000008,0x0004002b,0x0000000a,0x00000038,
+ 0x00000001,0x0006002c,0x0000000b,0x00000039,0x00000037,0x00000037,0x00000038,0x00050036,
+ 0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+ 0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+ 0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+ 0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+ 0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+ 0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+ 0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+ 0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+ 0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+ 0x0004003d,0x00000007,0x00000029,0x00000009,0x0007004f,0x00000010,0x0000002a,0x00000029,
+ 0x00000029,0x00000000,0x00000001,0x00050080,0x00000010,0x0000002b,0x00000028,0x0000002a,
+ 0x00050041,0x0000002d,0x0000002e,0x00000009,0x0000002c,0x0004003d,0x00000006,0x0000002f,
+ 0x0000002e,0x00050051,0x00000006,0x00000030,0x0000002b,0x00000000,0x00050051,0x00000006,
+ 0x00000031,0x0000002b,0x00000001,0x00060050,0x00000007,0x00000032,0x00000030,0x00000031,
+ 0x0000002f,0x00050041,0x00000034,0x00000035,0x00000017,0x00000033,0x0004003d,0x00000014,
+ 0x00000036,0x00000035,0x00040063,0x00000025,0x00000032,0x00000036,0x000200f9,0x00000021,
+ 0x000200f8,0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_array_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform uimage2DArray dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+ imageStore(dst, ivec3(u_info.dst_offset.xy + thread_id.xy, thread_id.z), u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000039,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+ 0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+ 0x00000019,0x00040047,0x00000038,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+ 0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+ 0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+ 0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+ 0x00000002,0x00090019,0x00000021,0x0000000a,0x00000001,0x00000000,0x00000001,0x00000000,
+ 0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+ 0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x0000000a,
+ 0x0000002b,0x00000002,0x00040020,0x0000002c,0x00000007,0x00000006,0x0004002b,0x00000006,
+ 0x00000032,0x00000000,0x00040020,0x00000033,0x00000009,0x00000013,0x0004002b,0x0000000a,
+ 0x00000036,0x00000008,0x0004002b,0x0000000a,0x00000037,0x00000001,0x0006002c,0x0000000b,
+ 0x00000038,0x00000036,0x00000036,0x00000037,0x00050036,0x00000002,0x00000004,0x00000000,
+ 0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+ 0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+ 0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+ 0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+ 0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+ 0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+ 0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+ 0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+ 0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x0004003d,0x00000007,0x00000028,
+ 0x00000009,0x0007004f,0x00000010,0x00000029,0x00000028,0x00000028,0x00000000,0x00000001,
+ 0x00050080,0x00000010,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+ 0x00000009,0x0000002b,0x0004003d,0x00000006,0x0000002e,0x0000002d,0x00050051,0x00000006,
+ 0x0000002f,0x0000002a,0x00000000,0x00050051,0x00000006,0x00000030,0x0000002a,0x00000001,
+ 0x00060050,0x00000007,0x00000031,0x0000002f,0x00000030,0x0000002e,0x00050041,0x00000033,
+ 0x00000034,0x00000016,0x00000032,0x0004003d,0x00000013,0x00000035,0x00000034,0x00040063,
+ 0x00000024,0x00000031,0x00000035,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,
+ 0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform image2D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+ imageStore(dst, u_info.dst_offset.xy + thread_id.xy, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000033,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+ 0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+ 0x00000019,0x00040047,0x00000032,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+ 0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+ 0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+ 0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+ 0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000001,
+ 0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+ 0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+ 0x00000001,0x0004002b,0x00000006,0x0000002c,0x00000000,0x00040020,0x0000002d,0x00000009,
+ 0x00000014,0x0004002b,0x0000000a,0x00000030,0x00000008,0x0004002b,0x0000000a,0x00000031,
+ 0x00000001,0x0006002c,0x0000000b,0x00000032,0x00000030,0x00000030,0x00000031,0x00050036,
+ 0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+ 0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+ 0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+ 0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+ 0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+ 0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+ 0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+ 0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+ 0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+ 0x0004003d,0x00000007,0x00000029,0x00000009,0x0007004f,0x00000010,0x0000002a,0x00000029,
+ 0x00000029,0x00000000,0x00000001,0x00050080,0x00000010,0x0000002b,0x00000028,0x0000002a,
+ 0x00050041,0x0000002d,0x0000002e,0x00000017,0x0000002c,0x0004003d,0x00000014,0x0000002f,
+ 0x0000002e,0x00040063,0x00000025,0x0000002b,0x0000002f,0x000200f9,0x00000021,0x000200f8,
+ 0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_2d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8) in;
+
+layout(binding = 0)
+writeonly uniform uimage2D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent.xy)))
+ imageStore(dst, u_info.dst_offset.xy + thread_id.xy, u_info.clear_value);
+}
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000032,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+ 0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+ 0x00000019,0x00040047,0x00000031,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+ 0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+ 0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+ 0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+ 0x00000002,0x00090019,0x00000021,0x0000000a,0x00000001,0x00000000,0x00000000,0x00000000,
+ 0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+ 0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,
+ 0x0000002b,0x00000000,0x00040020,0x0000002c,0x00000009,0x00000013,0x0004002b,0x0000000a,
+ 0x0000002f,0x00000008,0x0004002b,0x0000000a,0x00000030,0x00000001,0x0006002c,0x0000000b,
+ 0x00000031,0x0000002f,0x0000002f,0x00000030,0x00050036,0x00000002,0x00000004,0x00000000,
+ 0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+ 0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+ 0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+ 0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+ 0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+ 0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+ 0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+ 0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+ 0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x0004003d,0x00000007,0x00000028,
+ 0x00000009,0x0007004f,0x00000010,0x00000029,0x00000028,0x00000028,0x00000000,0x00000001,
+ 0x00050080,0x00000010,0x0000002a,0x00000027,0x00000029,0x00050041,0x0000002c,0x0000002d,
+ 0x00000016,0x0000002b,0x0004003d,0x00000013,0x0000002e,0x0000002d,0x00040063,0x00000024,
+ 0x0000002a,0x0000002e,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_3d_float_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
+
+layout(binding = 0)
+writeonly uniform image3D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ vec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent)))
+ imageStore(dst, ivec3(u_info.dst_offset.xy, 0) + thread_id.xyz, u_info.clear_value);
+}
+
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000035,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000015,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000015,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000015,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000015,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000017,0x6e695f75,0x00006f66,0x00030005,0x00000024,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000015,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000015,0x00000001,0x00000023,0x00000010,0x00050048,0x00000015,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000015,0x00000002,0x00040047,0x00000024,
+ 0x00000022,0x00000000,0x00040047,0x00000024,0x00000021,0x00000000,0x00030047,0x00000024,
+ 0x00000019,0x00040047,0x00000034,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00030016,0x00000013,0x00000020,0x00040017,0x00000014,0x00000013,
+ 0x00000004,0x0005001e,0x00000015,0x00000014,0x00000010,0x00000010,0x00040020,0x00000016,
+ 0x00000009,0x00000015,0x0004003b,0x00000016,0x00000017,0x00000009,0x0004002b,0x00000006,
+ 0x00000018,0x00000002,0x00040020,0x00000019,0x00000009,0x00000010,0x00020014,0x0000001c,
+ 0x00040017,0x0000001d,0x0000001c,0x00000002,0x00090019,0x00000022,0x00000013,0x00000002,
+ 0x00000000,0x00000000,0x00000000,0x00000002,0x00000000,0x00040020,0x00000023,0x00000000,
+ 0x00000022,0x0004003b,0x00000023,0x00000024,0x00000000,0x0004002b,0x00000006,0x00000026,
+ 0x00000001,0x0004002b,0x00000006,0x00000029,0x00000000,0x00040020,0x0000002f,0x00000009,
+ 0x00000014,0x0004002b,0x0000000a,0x00000032,0x00000008,0x0004002b,0x0000000a,0x00000033,
+ 0x00000001,0x0006002c,0x0000000b,0x00000034,0x00000032,0x00000032,0x00000033,0x00050036,
+ 0x00000002,0x00000004,0x00000000,0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,
+ 0x00000009,0x00000007,0x0004003d,0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,
+ 0x0000000f,0x0000000e,0x0003003e,0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,
+ 0x00000009,0x0007004f,0x00000010,0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,
+ 0x00050041,0x00000019,0x0000001a,0x00000017,0x00000018,0x0004003d,0x00000010,0x0000001b,
+ 0x0000001a,0x000500b1,0x0000001d,0x0000001e,0x00000012,0x0000001b,0x0004009b,0x0000001c,
+ 0x0000001f,0x0000001e,0x000300f7,0x00000021,0x00000000,0x000400fa,0x0000001f,0x00000020,
+ 0x00000021,0x000200f8,0x00000020,0x0004003d,0x00000022,0x00000025,0x00000024,0x00050041,
+ 0x00000019,0x00000027,0x00000017,0x00000026,0x0004003d,0x00000010,0x00000028,0x00000027,
+ 0x00050051,0x00000006,0x0000002a,0x00000028,0x00000000,0x00050051,0x00000006,0x0000002b,
+ 0x00000028,0x00000001,0x00060050,0x00000007,0x0000002c,0x0000002a,0x0000002b,0x00000029,
+ 0x0004003d,0x00000007,0x0000002d,0x00000009,0x00050080,0x00000007,0x0000002e,0x0000002c,
+ 0x0000002d,0x00050041,0x0000002f,0x00000030,0x00000017,0x00000029,0x0004003d,0x00000014,
+ 0x00000031,0x00000030,0x00040063,0x00000025,0x0000002e,0x00000031,0x000200f9,0x00000021,
+ 0x000200f8,0x00000021,0x000100fd,0x00010038
+};
+
+const uint32_t cs_clear_uav_image_3d_uint_spv[] = {
+#if 0
+#version 450
+
+layout(local_size_x = 8, local_size_y = 8, local_size_z = 1) in;
+
+layout(binding = 0)
+writeonly uniform uimage3D dst;
+
+layout(push_constant)
+uniform u_info_t {
+ uvec4 clear_value;
+ ivec2 dst_offset;
+ ivec2 dst_extent;
+} u_info;
+
+void main() {
+ ivec3 thread_id = ivec3(gl_GlobalInvocationID);
+
+ if (all(lessThan(thread_id.xy, u_info.dst_extent)))
+ imageStore(dst, ivec3(u_info.dst_offset.xy, 0) + thread_id.xyz, u_info.clear_value);
+}
+
+#endif
+ 0x07230203,0x00010000,0x00080007,0x00000034,0x00000000,0x00020011,0x00000001,0x00020011,
+ 0x00000038,0x0006000b,0x00000001,0x4c534c47,0x6474732e,0x3035342e,0x00000000,0x0003000e,
+ 0x00000000,0x00000001,0x0006000f,0x00000005,0x00000004,0x6e69616d,0x00000000,0x0000000d,
+ 0x00060010,0x00000004,0x00000011,0x00000008,0x00000008,0x00000001,0x00030003,0x00000002,
+ 0x000001c2,0x00040005,0x00000004,0x6e69616d,0x00000000,0x00050005,0x00000009,0x65726874,
+ 0x695f6461,0x00000064,0x00080005,0x0000000d,0x475f6c67,0x61626f6c,0x766e496c,0x7461636f,
+ 0x496e6f69,0x00000044,0x00050005,0x00000014,0x6e695f75,0x745f6f66,0x00000000,0x00060006,
+ 0x00000014,0x00000000,0x61656c63,0x61765f72,0x0065756c,0x00060006,0x00000014,0x00000001,
+ 0x5f747364,0x7366666f,0x00007465,0x00060006,0x00000014,0x00000002,0x5f747364,0x65747865,
+ 0x0000746e,0x00040005,0x00000016,0x6e695f75,0x00006f66,0x00030005,0x00000023,0x00747364,
+ 0x00040047,0x0000000d,0x0000000b,0x0000001c,0x00050048,0x00000014,0x00000000,0x00000023,
+ 0x00000000,0x00050048,0x00000014,0x00000001,0x00000023,0x00000010,0x00050048,0x00000014,
+ 0x00000002,0x00000023,0x00000018,0x00030047,0x00000014,0x00000002,0x00040047,0x00000023,
+ 0x00000022,0x00000000,0x00040047,0x00000023,0x00000021,0x00000000,0x00030047,0x00000023,
+ 0x00000019,0x00040047,0x00000033,0x0000000b,0x00000019,0x00020013,0x00000002,0x00030021,
+ 0x00000003,0x00000002,0x00040015,0x00000006,0x00000020,0x00000001,0x00040017,0x00000007,
+ 0x00000006,0x00000003,0x00040020,0x00000008,0x00000007,0x00000007,0x00040015,0x0000000a,
+ 0x00000020,0x00000000,0x00040017,0x0000000b,0x0000000a,0x00000003,0x00040020,0x0000000c,
+ 0x00000001,0x0000000b,0x0004003b,0x0000000c,0x0000000d,0x00000001,0x00040017,0x00000010,
+ 0x00000006,0x00000002,0x00040017,0x00000013,0x0000000a,0x00000004,0x0005001e,0x00000014,
+ 0x00000013,0x00000010,0x00000010,0x00040020,0x00000015,0x00000009,0x00000014,0x0004003b,
+ 0x00000015,0x00000016,0x00000009,0x0004002b,0x00000006,0x00000017,0x00000002,0x00040020,
+ 0x00000018,0x00000009,0x00000010,0x00020014,0x0000001b,0x00040017,0x0000001c,0x0000001b,
+ 0x00000002,0x00090019,0x00000021,0x0000000a,0x00000002,0x00000000,0x00000000,0x00000000,
+ 0x00000002,0x00000000,0x00040020,0x00000022,0x00000000,0x00000021,0x0004003b,0x00000022,
+ 0x00000023,0x00000000,0x0004002b,0x00000006,0x00000025,0x00000001,0x0004002b,0x00000006,
+ 0x00000028,0x00000000,0x00040020,0x0000002e,0x00000009,0x00000013,0x0004002b,0x0000000a,
+ 0x00000031,0x00000008,0x0004002b,0x0000000a,0x00000032,0x00000001,0x0006002c,0x0000000b,
+ 0x00000033,0x00000031,0x00000031,0x00000032,0x00050036,0x00000002,0x00000004,0x00000000,
+ 0x00000003,0x000200f8,0x00000005,0x0004003b,0x00000008,0x00000009,0x00000007,0x0004003d,
+ 0x0000000b,0x0000000e,0x0000000d,0x0004007c,0x00000007,0x0000000f,0x0000000e,0x0003003e,
+ 0x00000009,0x0000000f,0x0004003d,0x00000007,0x00000011,0x00000009,0x0007004f,0x00000010,
+ 0x00000012,0x00000011,0x00000011,0x00000000,0x00000001,0x00050041,0x00000018,0x00000019,
+ 0x00000016,0x00000017,0x0004003d,0x00000010,0x0000001a,0x00000019,0x000500b1,0x0000001c,
+ 0x0000001d,0x00000012,0x0000001a,0x0004009b,0x0000001b,0x0000001e,0x0000001d,0x000300f7,
+ 0x00000020,0x00000000,0x000400fa,0x0000001e,0x0000001f,0x00000020,0x000200f8,0x0000001f,
+ 0x0004003d,0x00000021,0x00000024,0x00000023,0x00050041,0x00000018,0x00000026,0x00000016,
+ 0x00000025,0x0004003d,0x00000010,0x00000027,0x00000026,0x00050051,0x00000006,0x00000029,
+ 0x00000027,0x00000000,0x00050051,0x00000006,0x0000002a,0x00000027,0x00000001,0x00060050,
+ 0x00000007,0x0000002b,0x00000029,0x0000002a,0x00000028,0x0004003d,0x00000007,0x0000002c,
+ 0x00000009,0x00050080,0x00000007,0x0000002d,0x0000002b,0x0000002c,0x00050041,0x0000002e,
+ 0x0000002f,0x00000016,0x00000028,0x0004003d,0x00000013,0x00000030,0x0000002f,0x00040063,
+ 0x00000024,0x0000002d,0x00000030,0x000200f9,0x00000020,0x000200f8,0x00000020,0x000100fd,
+ 0x00010038
+};
+
+#endif /* __VKD3D_SPV_SHADERS_H */
--
2.24.0
More information about the wine-devel
mailing list