Signed-off-by: Philip Rebohle philip.rebohle@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 */