vkd3d: Restore old ClearUnorderedAccessView implementation.
The current code uses D3D12 abstractions to create pipelines but issues raw Vulkan API calls to actually implement the functionality, which means the code makes assumptions about the exact descriptor set layout and push constant layout, which is generally a bad idea now that we have multiple code paths for root constants etc. Signed-off-by: Philip Rebohle <philip.rebohle@tu-dortmund.de>
This commit is contained in:
Philip Rebohle
Hans-Kristian Arntzen
parent
edc614ab2e
commit
957aec4785
|
|
@ -108,6 +108,7 @@ 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 \
|
||||
|
|
|
|||
|
|
@ -4993,81 +4993,18 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearRenderTargetView(d3d12_com
|
|||
&clear_value, rect_count, rects);
|
||||
}
|
||||
|
||||
struct vkd3d_uav_clear_pipeline
|
||||
{
|
||||
VkDescriptorSetLayout vk_set_layout;
|
||||
VkPipelineLayout vk_pipeline_layout;
|
||||
VkPipeline vk_pipeline;
|
||||
VkExtent3D group_size;
|
||||
};
|
||||
|
||||
static void vkd3d_uav_clear_state_get_buffer_pipeline(const struct vkd3d_uav_clear_state *state,
|
||||
enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
|
||||
{
|
||||
const struct vkd3d_uav_clear_pipelines *pipelines;
|
||||
|
||||
pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
|
||||
info->vk_set_layout = state->vk_set_layout_buffer;
|
||||
info->vk_pipeline_layout = state->vk_pipeline_layout_buffer;
|
||||
info->vk_pipeline = pipelines->buffer;
|
||||
info->group_size = (VkExtent3D){128, 1, 1};
|
||||
}
|
||||
|
||||
static void vkd3d_uav_clear_state_get_image_pipeline(const struct vkd3d_uav_clear_state *state,
|
||||
VkImageViewType image_view_type, enum vkd3d_format_type format_type, struct vkd3d_uav_clear_pipeline *info)
|
||||
{
|
||||
const struct vkd3d_uav_clear_pipelines *pipelines;
|
||||
|
||||
pipelines = format_type == VKD3D_FORMAT_TYPE_UINT ? &state->pipelines_uint : &state->pipelines_float;
|
||||
info->vk_set_layout = state->vk_set_layout_image;
|
||||
info->vk_pipeline_layout = state->vk_pipeline_layout_image;
|
||||
|
||||
switch (image_view_type)
|
||||
{
|
||||
case VK_IMAGE_VIEW_TYPE_1D:
|
||||
info->vk_pipeline = pipelines->image_1d;
|
||||
info->group_size = (VkExtent3D){64, 1, 1};
|
||||
break;
|
||||
|
||||
case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
|
||||
info->vk_pipeline = pipelines->image_1d_array;
|
||||
info->group_size = (VkExtent3D){64, 1, 1};
|
||||
break;
|
||||
|
||||
case VK_IMAGE_VIEW_TYPE_2D:
|
||||
info->vk_pipeline = pipelines->image_2d;
|
||||
info->group_size = (VkExtent3D){8, 8, 1};
|
||||
break;
|
||||
|
||||
case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
|
||||
info->vk_pipeline = pipelines->image_2d_array;
|
||||
info->group_size = (VkExtent3D){8, 8, 1};
|
||||
break;
|
||||
|
||||
case VK_IMAGE_VIEW_TYPE_3D:
|
||||
info->vk_pipeline = pipelines->image_3d;
|
||||
info->group_size = (VkExtent3D){8, 8, 1};
|
||||
break;
|
||||
|
||||
default:
|
||||
ERR("Unhandled view type %#x.\n", image_view_type);
|
||||
info->vk_pipeline = VK_NULL_HANDLE;
|
||||
info->group_size = (VkExtent3D){0, 0, 0};
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
|
||||
struct d3d12_resource *resource, struct vkd3d_view *view, const VkClearColorValue *clear_colour,
|
||||
unsigned int rect_count, const D3D12_RECT *rects)
|
||||
struct d3d12_resource *resource, struct vkd3d_view *view, const VkClearColorValue *clear_color,
|
||||
UINT rect_count, const D3D12_RECT *rects)
|
||||
{
|
||||
const struct vkd3d_vk_device_procs *vk_procs = &list->device->vk_procs;
|
||||
unsigned int i, miplevel_idx, layer_count;
|
||||
struct vkd3d_uav_clear_pipeline pipeline;
|
||||
struct vkd3d_uav_clear_args clear_args;
|
||||
struct vkd3d_clear_uav_pipeline pipeline;
|
||||
struct vkd3d_clear_uav_args clear_args;
|
||||
VkDescriptorImageInfo image_info;
|
||||
D3D12_RECT full_rect, curr_rect;
|
||||
VkWriteDescriptorSet write_set;
|
||||
VkExtent3D workgroup_size;
|
||||
|
||||
d3d12_command_list_track_resource_usage(list, resource);
|
||||
d3d12_command_list_end_current_render_pass(list);
|
||||
|
|
@ -5078,7 +5015,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
|
|||
if (!d3d12_command_allocator_add_view(list->allocator, view))
|
||||
WARN("Failed to add view.\n");
|
||||
|
||||
clear_args.colour = *clear_colour;
|
||||
clear_args.clear_color = *clear_color;
|
||||
|
||||
write_set.sType = VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET;
|
||||
write_set.pNext = NULL;
|
||||
|
|
@ -5086,19 +5023,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
|
|||
write_set.dstArrayElement = 0;
|
||||
write_set.descriptorCount = 1;
|
||||
|
||||
if (d3d12_resource_is_buffer(resource))
|
||||
{
|
||||
write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER;
|
||||
write_set.pImageInfo = NULL;
|
||||
write_set.pBufferInfo = NULL;
|
||||
write_set.pTexelBufferView = &view->u.vk_buffer_view;
|
||||
|
||||
miplevel_idx = 0;
|
||||
layer_count = 1;
|
||||
vkd3d_uav_clear_state_get_buffer_pipeline(&list->device->uav_clear_state,
|
||||
view->format->type, &pipeline);
|
||||
}
|
||||
else
|
||||
if (d3d12_resource_is_texture(resource))
|
||||
{
|
||||
image_info.sampler = VK_NULL_HANDLE;
|
||||
image_info.imageView = view->u.vk_image_view;
|
||||
|
|
@ -5113,8 +5038,24 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
|
|||
layer_count = view->info.texture.vk_view_type == VK_IMAGE_VIEW_TYPE_3D
|
||||
? d3d12_resource_desc_get_depth(&resource->desc, miplevel_idx)
|
||||
: view->info.texture.layer_count;
|
||||
vkd3d_uav_clear_state_get_image_pipeline(&list->device->uav_clear_state,
|
||||
view->info.texture.vk_view_type, view->format->type, &pipeline);
|
||||
pipeline = vkd3d_clear_uav_ops_get_clear_image_pipeline(
|
||||
&list->device->meta_ops.clear_uav, view->info.texture.vk_view_type,
|
||||
view->format->type == VKD3D_FORMAT_TYPE_UINT);
|
||||
workgroup_size = vkd3d_get_clear_image_uav_workgroup_size(view->info.texture.vk_view_type);
|
||||
}
|
||||
else
|
||||
{
|
||||
write_set.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER;
|
||||
write_set.pImageInfo = NULL;
|
||||
write_set.pBufferInfo = NULL;
|
||||
write_set.pTexelBufferView = &view->u.vk_buffer_view;
|
||||
|
||||
miplevel_idx = 0;
|
||||
layer_count = 1;
|
||||
pipeline = vkd3d_clear_uav_ops_get_clear_buffer_pipeline(
|
||||
&list->device->meta_ops.clear_uav,
|
||||
view->format->type == VKD3D_FORMAT_TYPE_UINT);
|
||||
workgroup_size = vkd3d_get_clear_buffer_uav_workgroup_size();
|
||||
}
|
||||
|
||||
if (!(write_set.dstSet = d3d12_command_allocator_allocate_descriptor_set(
|
||||
|
|
@ -5131,52 +5072,55 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
|
|||
full_rect.top = 0;
|
||||
full_rect.bottom = d3d12_resource_desc_get_height(&resource->desc, miplevel_idx);
|
||||
|
||||
if (!rect_count)
|
||||
/* clear full resource if no rects are specified */
|
||||
curr_rect = full_rect;
|
||||
|
||||
VK_CALL(vkCmdBindPipeline(list->vk_command_buffer,
|
||||
VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline));
|
||||
|
||||
VK_CALL(vkCmdBindDescriptorSets(list->vk_command_buffer,
|
||||
VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline_layout,
|
||||
0, 1, &write_set.dstSet, 0, NULL));
|
||||
|
||||
for (i = 0; i < rect_count || !i; i++)
|
||||
{
|
||||
rects = &full_rect;
|
||||
rect_count = 1;
|
||||
}
|
||||
if (rect_count)
|
||||
{
|
||||
/* clamp to actual resource region and skip empty rects */
|
||||
curr_rect.left = max(rects[i].left, full_rect.left);
|
||||
curr_rect.top = max(rects[i].top, full_rect.top);
|
||||
curr_rect.right = min(rects[i].right, full_rect.right);
|
||||
curr_rect.bottom = min(rects[i].bottom, full_rect.bottom);
|
||||
|
||||
VK_CALL(vkCmdBindPipeline(list->vk_command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE, pipeline.vk_pipeline));
|
||||
|
||||
VK_CALL(vkCmdBindDescriptorSets(list->vk_command_buffer, VK_PIPELINE_BIND_POINT_COMPUTE,
|
||||
pipeline.vk_pipeline_layout, 0, 1, &write_set.dstSet, 0, NULL));
|
||||
|
||||
for (i = 0; i < rect_count; ++i)
|
||||
{
|
||||
/* Clamp to the actual resource region and skip empty rectangles. */
|
||||
curr_rect.left = max(rects[i].left, full_rect.left);
|
||||
curr_rect.top = max(rects[i].top, full_rect.top);
|
||||
curr_rect.right = min(rects[i].right, full_rect.right);
|
||||
curr_rect.bottom = min(rects[i].bottom, full_rect.bottom);
|
||||
|
||||
if (curr_rect.left >= curr_rect.right || curr_rect.top >= curr_rect.bottom)
|
||||
continue;
|
||||
if (curr_rect.left >= curr_rect.right || curr_rect.top >= curr_rect.bottom)
|
||||
continue;
|
||||
}
|
||||
|
||||
clear_args.offset.x = curr_rect.left;
|
||||
clear_args.offset.y = curr_rect.top;
|
||||
clear_args.extent.width = curr_rect.right - curr_rect.left;
|
||||
clear_args.extent.height = curr_rect.bottom - curr_rect.top;
|
||||
|
||||
VK_CALL(vkCmdPushConstants(list->vk_command_buffer, pipeline.vk_pipeline_layout,
|
||||
VK_SHADER_STAGE_COMPUTE_BIT, 0, sizeof(clear_args), &clear_args));
|
||||
VK_CALL(vkCmdPushConstants(list->vk_command_buffer,
|
||||
pipeline.vk_pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT,
|
||||
0, sizeof(clear_args), &clear_args));
|
||||
|
||||
VK_CALL(vkCmdDispatch(list->vk_command_buffer,
|
||||
vkd3d_compute_workgroup_count(clear_args.extent.width, pipeline.group_size.width),
|
||||
vkd3d_compute_workgroup_count(clear_args.extent.height, pipeline.group_size.height),
|
||||
vkd3d_compute_workgroup_count(layer_count, pipeline.group_size.depth)));
|
||||
vkd3d_compute_workgroup_count(clear_args.extent.width, workgroup_size.width),
|
||||
vkd3d_compute_workgroup_count(clear_args.extent.height, workgroup_size.height),
|
||||
vkd3d_compute_workgroup_count(layer_count, workgroup_size.depth)));
|
||||
}
|
||||
}
|
||||
|
||||
static const struct vkd3d_format *vkd3d_fixup_clear_uav_uint_colour(struct d3d12_device *device,
|
||||
DXGI_FORMAT dxgi_format, VkClearColorValue *colour)
|
||||
static const struct vkd3d_format *vkd3d_fixup_clear_uav_uint_color(struct d3d12_device *device,
|
||||
DXGI_FORMAT dxgi_format, VkClearColorValue *color)
|
||||
{
|
||||
switch (dxgi_format)
|
||||
{
|
||||
case DXGI_FORMAT_R11G11B10_FLOAT:
|
||||
colour->uint32[0] = (colour->uint32[0] & 0x7ff)
|
||||
| ((colour->uint32[1] & 0x7ff) << 11)
|
||||
| ((colour->uint32[2] & 0x3ff) << 22);
|
||||
color->uint32[0] = (color->uint32[0] & 0x7FF)
|
||||
| ((color->uint32[1] & 0x7FF) << 11)
|
||||
| ((color->uint32[2] & 0x3FF) << 22);
|
||||
return vkd3d_get_format(device, DXGI_FORMAT_R32_UINT, false);
|
||||
|
||||
default:
|
||||
|
|
@ -5189,61 +5133,66 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewUint(d3
|
|||
const UINT values[4], UINT rect_count, const D3D12_RECT *rects)
|
||||
{
|
||||
struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList(iface);
|
||||
struct d3d12_device *device = list->device;
|
||||
struct vkd3d_view *view, *uint_view = NULL;
|
||||
struct vkd3d_view *base_view, *uint_view;
|
||||
struct vkd3d_texture_view_desc view_desc;
|
||||
const struct vkd3d_format *uint_format;
|
||||
struct d3d12_resource *resource_impl;
|
||||
VkClearColorValue colour;
|
||||
VkClearColorValue color;
|
||||
|
||||
TRACE("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p.\n",
|
||||
iface, gpu_handle.ptr, cpu_handle.ptr, resource, values, rect_count, rects);
|
||||
|
||||
resource_impl = unsafe_impl_from_ID3D12Resource(resource);
|
||||
view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
|
||||
memcpy(colour.uint32, values, sizeof(colour.uint32));
|
||||
memcpy(color.uint32, values, sizeof(color.uint32));
|
||||
|
||||
if (view->format->type != VKD3D_FORMAT_TYPE_UINT)
|
||||
resource_impl = unsafe_impl_from_ID3D12Resource(resource);
|
||||
|
||||
base_view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
|
||||
uint_view = NULL;
|
||||
|
||||
if (base_view->format->type != VKD3D_FORMAT_TYPE_UINT)
|
||||
{
|
||||
if (!(uint_format = vkd3d_find_uint_format(device, view->format->dxgi_format))
|
||||
&& !(uint_format = vkd3d_fixup_clear_uav_uint_colour(device, view->format->dxgi_format, &colour)))
|
||||
uint_format = vkd3d_find_uint_format(list->device, base_view->format->dxgi_format);
|
||||
|
||||
if (!uint_format && !(uint_format = vkd3d_fixup_clear_uav_uint_color(
|
||||
list->device, base_view->format->dxgi_format, &color)))
|
||||
{
|
||||
ERR("Unhandled format %#x.\n", view->format->dxgi_format);
|
||||
ERR("Unhandled format %d.\n", base_view->format->dxgi_format);
|
||||
return;
|
||||
}
|
||||
|
||||
if (d3d12_resource_is_buffer(resource_impl))
|
||||
{
|
||||
if (!vkd3d_create_buffer_view(device, resource_impl->u.vk_buffer, uint_format,
|
||||
view->info.buffer.offset, view->info.buffer.size, &uint_view))
|
||||
{
|
||||
ERR("Failed to create buffer view.\n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
else
|
||||
if (d3d12_resource_is_texture(resource_impl))
|
||||
{
|
||||
memset(&view_desc, 0, sizeof(view_desc));
|
||||
view_desc.view_type = view->info.texture.vk_view_type;
|
||||
view_desc.view_type = base_view->info.texture.vk_view_type;
|
||||
view_desc.format = uint_format;
|
||||
view_desc.miplevel_idx = view->info.texture.miplevel_idx;
|
||||
view_desc.miplevel_idx = base_view->info.texture.miplevel_idx;
|
||||
view_desc.miplevel_count = 1;
|
||||
view_desc.layer_idx = view->info.texture.layer_idx;
|
||||
view_desc.layer_count = view->info.texture.layer_count;
|
||||
view_desc.layer_idx = base_view->info.texture.layer_idx;
|
||||
view_desc.layer_count = base_view->info.texture.layer_count;
|
||||
view_desc.allowed_swizzle = false;
|
||||
|
||||
if (!vkd3d_create_texture_view(device, resource_impl->u.vk_image, &view_desc, &uint_view))
|
||||
if (!vkd3d_create_texture_view(list->device, resource_impl->u.vk_image, &view_desc, &uint_view))
|
||||
{
|
||||
ERR("Failed to create image view.\n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
view = uint_view;
|
||||
else
|
||||
{
|
||||
if (!vkd3d_create_buffer_view(list->device, resource_impl->u.vk_buffer, uint_format,
|
||||
base_view->info.buffer.offset, base_view->info.buffer.size, &uint_view))
|
||||
{
|
||||
ERR("Failed to create buffer view.\n");
|
||||
return;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
d3d12_command_list_clear_uav(list, resource_impl, view, &colour, rect_count, rects);
|
||||
d3d12_command_list_clear_uav(list, resource_impl,
|
||||
uint_view ? uint_view : base_view, &color, rect_count, rects);
|
||||
|
||||
if (uint_view)
|
||||
vkd3d_view_decref(uint_view, device);
|
||||
vkd3d_view_decref(uint_view, list->device);
|
||||
}
|
||||
|
||||
static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(d3d12_command_list_iface *iface,
|
||||
|
|
@ -5252,17 +5201,18 @@ static void STDMETHODCALLTYPE d3d12_command_list_ClearUnorderedAccessViewFloat(d
|
|||
{
|
||||
struct d3d12_command_list *list = impl_from_ID3D12GraphicsCommandList(iface);
|
||||
struct d3d12_resource *resource_impl;
|
||||
VkClearColorValue colour;
|
||||
struct vkd3d_view *view;
|
||||
VkClearColorValue color;
|
||||
|
||||
TRACE("iface %p, gpu_handle %#"PRIx64", cpu_handle %lx, resource %p, values %p, rect_count %u, rects %p.\n",
|
||||
iface, gpu_handle.ptr, cpu_handle.ptr, resource, values, rect_count, rects);
|
||||
|
||||
memcpy(color.float32, values, sizeof(color.float32));
|
||||
|
||||
resource_impl = unsafe_impl_from_ID3D12Resource(resource);
|
||||
view = d3d12_desc_from_cpu_handle(cpu_handle)->u.view;
|
||||
memcpy(colour.float32, values, sizeof(colour.float32));
|
||||
|
||||
d3d12_command_list_clear_uav(list, resource_impl, view, &colour, rect_count, rects);
|
||||
d3d12_command_list_clear_uav(list, resource_impl, view, &color, rect_count, rects);
|
||||
}
|
||||
|
||||
static void STDMETHODCALLTYPE d3d12_command_list_DiscardResource(d3d12_command_list_iface *iface,
|
||||
|
|
|
|||
|
|
@ -2184,7 +2184,7 @@ static void d3d12_device_destroy(struct d3d12_device *device)
|
|||
vkd3d_private_store_destroy(&device->private_store);
|
||||
|
||||
vkd3d_cleanup_format_info(device);
|
||||
vkd3d_uav_clear_state_cleanup(&device->uav_clear_state, device);
|
||||
vkd3d_meta_ops_cleanup(&device->meta_ops, device);
|
||||
vkd3d_bindless_state_cleanup(&device->bindless_state, device);
|
||||
vkd3d_destroy_null_resources(&device->null_resources, device);
|
||||
vkd3d_gpu_va_allocator_cleanup(&device->gpu_va_allocator);
|
||||
|
|
@ -4465,7 +4465,7 @@ static HRESULT d3d12_device_init(struct d3d12_device *device,
|
|||
if (FAILED(hr = vkd3d_bindless_state_init(&device->bindless_state, device)))
|
||||
goto out_destroy_null_resources;
|
||||
|
||||
if (FAILED(hr = vkd3d_uav_clear_state_init(&device->uav_clear_state, device)))
|
||||
if (FAILED(hr = vkd3d_meta_ops_init(&device->meta_ops, device)))
|
||||
goto out_cleanup_bindless_state;
|
||||
|
||||
vkd3d_render_pass_cache_init(&device->render_pass_cache);
|
||||
|
|
|
|||
|
|
@ -0,0 +1,343 @@
|
|||
/*
|
||||
* Copyright 2019 Philip Rebohle
|
||||
*
|
||||
* This library is free software; you can redistribute it and/or
|
||||
* modify it under the terms of the GNU Lesser General Public
|
||||
* License as published by the Free Software Foundation; either
|
||||
* version 2.1 of the License, or (at your option) any later version.
|
||||
*
|
||||
* This library is distributed in the hope that it will be useful,
|
||||
* but WITHOUT ANY WARRANTY; without even the implied warranty of
|
||||
* MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE. See the GNU
|
||||
* Lesser General Public License for more details.
|
||||
*
|
||||
* You should have received a copy of the GNU Lesser General Public
|
||||
* License along with this library; if not, write to the Free Software
|
||||
* Foundation, Inc., 51 Franklin St, Fifth Floor, Boston, MA 02110-1301, USA
|
||||
*/
|
||||
|
||||
#include "vkd3d_private.h"
|
||||
#include "vkd3d_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.buffer,
|
||||
&meta_clear_uav_ops->vk_pipeline_layout_buffer,
|
||||
SPIRV_CODE(cs_clear_uav_buffer_float_spv) },
|
||||
{ &meta_clear_uav_ops->clear_float.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.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.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.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.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.buffer,
|
||||
&meta_clear_uav_ops->vk_pipeline_layout_buffer,
|
||||
SPIRV_CODE(cs_clear_uav_buffer_uint_spv) },
|
||||
{ &meta_clear_uav_ops->clear_uint.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.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.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.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.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_cleanup(meta_clear_uav_ops, device);
|
||||
return hresult_from_vk_result(vr);
|
||||
}
|
||||
|
||||
void vkd3d_clear_uav_ops_cleanup(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]->buffer, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->image_1d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->image_2d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->image_3d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->image_1d_array, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipeline_sets[i]->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->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->image_1d;
|
||||
break;
|
||||
case VK_IMAGE_VIEW_TYPE_2D:
|
||||
info.vk_pipeline = pipelines->image_2d;
|
||||
break;
|
||||
case VK_IMAGE_VIEW_TYPE_3D:
|
||||
info.vk_pipeline = pipelines->image_3d;
|
||||
break;
|
||||
case VK_IMAGE_VIEW_TYPE_1D_ARRAY:
|
||||
info.vk_pipeline = pipelines->image_1d_array;
|
||||
break;
|
||||
case VK_IMAGE_VIEW_TYPE_2D_ARRAY:
|
||||
info.vk_pipeline = pipelines->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_cleanup(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device)
|
||||
{
|
||||
vkd3d_clear_uav_ops_cleanup(&meta_ops->clear_uav, device);
|
||||
return S_OK;
|
||||
}
|
||||
|
|
@ -18,7 +18,6 @@
|
|||
*/
|
||||
|
||||
#include "vkd3d_private.h"
|
||||
#include "vkd3d_shaders.h"
|
||||
|
||||
/* ID3D12RootSignature */
|
||||
static inline struct d3d12_root_signature *impl_from_ID3D12RootSignature(ID3D12RootSignature *iface)
|
||||
|
|
@ -3386,167 +3385,3 @@ HRESULT d3d12_pipeline_library_create(struct d3d12_device *device, const void *b
|
|||
*pipeline_library = object;
|
||||
return S_OK;
|
||||
}
|
||||
|
||||
static void vkd3d_uav_clear_pipelines_cleanup(struct vkd3d_uav_clear_pipelines *pipelines,
|
||||
struct d3d12_device *device)
|
||||
{
|
||||
const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
|
||||
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_3d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d_array, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_2d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d_array, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->image_1d, NULL));
|
||||
VK_CALL(vkDestroyPipeline(device->vk_device, pipelines->buffer, NULL));
|
||||
}
|
||||
|
||||
void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
|
||||
{
|
||||
const struct vkd3d_vk_device_procs *vk_procs = &device->vk_procs;
|
||||
|
||||
vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_uint, device);
|
||||
vkd3d_uav_clear_pipelines_cleanup(&state->pipelines_float, device);
|
||||
|
||||
VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_image, NULL));
|
||||
VK_CALL(vkDestroyPipelineLayout(device->vk_device, state->vk_pipeline_layout_buffer, NULL));
|
||||
|
||||
VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_image, NULL));
|
||||
VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, state->vk_set_layout_buffer, NULL));
|
||||
}
|
||||
|
||||
HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device)
|
||||
{
|
||||
struct vkd3d_shader_push_constant_buffer push_constant;
|
||||
struct vkd3d_shader_interface_info shader_interface;
|
||||
struct vkd3d_shader_resource_binding binding;
|
||||
VkDescriptorSetLayoutBinding set_binding;
|
||||
VkPushConstantRange push_constant_range;
|
||||
unsigned int i;
|
||||
HRESULT hr;
|
||||
|
||||
const struct
|
||||
{
|
||||
VkDescriptorSetLayout *set_layout;
|
||||
VkPipelineLayout *pipeline_layout;
|
||||
VkDescriptorType descriptor_type;
|
||||
}
|
||||
set_layouts[] =
|
||||
{
|
||||
{&state->vk_set_layout_buffer, &state->vk_pipeline_layout_buffer, VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER},
|
||||
{&state->vk_set_layout_image, &state->vk_pipeline_layout_image, VK_DESCRIPTOR_TYPE_STORAGE_IMAGE},
|
||||
};
|
||||
|
||||
const struct
|
||||
{
|
||||
VkPipeline *pipeline;
|
||||
VkPipelineLayout *pipeline_layout;
|
||||
D3D12_SHADER_BYTECODE code;
|
||||
}
|
||||
pipelines[] =
|
||||
{
|
||||
#define SHADER_CODE(name) {name, sizeof(name)}
|
||||
{&state->pipelines_float.buffer, &state->vk_pipeline_layout_buffer,
|
||||
SHADER_CODE(cs_uav_clear_buffer_float_code)},
|
||||
{&state->pipelines_float.image_1d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_1d_float_code)},
|
||||
{&state->pipelines_float.image_1d_array, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_1d_array_float_code)},
|
||||
{&state->pipelines_float.image_2d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_2d_float_code)},
|
||||
{&state->pipelines_float.image_2d_array, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_2d_array_float_code)},
|
||||
{&state->pipelines_float.image_3d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_3d_float_code)},
|
||||
|
||||
{&state->pipelines_uint.buffer, &state->vk_pipeline_layout_buffer,
|
||||
SHADER_CODE(cs_uav_clear_buffer_uint_code)},
|
||||
{&state->pipelines_uint.image_1d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_1d_uint_code)},
|
||||
{&state->pipelines_uint.image_1d_array, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_1d_array_uint_code)},
|
||||
{&state->pipelines_uint.image_2d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_2d_uint_code)},
|
||||
{&state->pipelines_uint.image_2d_array, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_2d_array_uint_code)},
|
||||
{&state->pipelines_uint.image_3d, &state->vk_pipeline_layout_image,
|
||||
SHADER_CODE(cs_uav_clear_3d_uint_code)},
|
||||
#undef SHADER_CODE
|
||||
};
|
||||
|
||||
memset(state, 0, sizeof(*state));
|
||||
|
||||
set_binding.binding = 0;
|
||||
set_binding.descriptorCount = 1;
|
||||
set_binding.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
|
||||
set_binding.pImmutableSamplers = NULL;
|
||||
|
||||
binding.type = VKD3D_SHADER_DESCRIPTOR_TYPE_UAV;
|
||||
binding.register_index = 0;
|
||||
binding.register_space = 0;
|
||||
binding.register_count = 1;
|
||||
binding.descriptor_table = 0;
|
||||
binding.descriptor_offset = 0;
|
||||
binding.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
|
||||
binding.binding.set = 0;
|
||||
binding.binding.binding = 0;
|
||||
|
||||
push_constant_range.stageFlags = VK_SHADER_STAGE_COMPUTE_BIT;
|
||||
push_constant_range.offset = 0;
|
||||
push_constant_range.size = sizeof(struct vkd3d_uav_clear_args);
|
||||
|
||||
push_constant.register_space = 0;
|
||||
push_constant.register_index = 0;
|
||||
push_constant.shader_visibility = VKD3D_SHADER_VISIBILITY_COMPUTE;
|
||||
push_constant.offset = 0;
|
||||
push_constant.size = sizeof(struct vkd3d_uav_clear_args);
|
||||
|
||||
for (i = 0; i < ARRAY_SIZE(set_layouts); ++i)
|
||||
{
|
||||
set_binding.descriptorType = set_layouts[i].descriptor_type;
|
||||
|
||||
if (FAILED(hr = vkd3d_create_descriptor_set_layout(device, 0, 1, &set_binding, set_layouts[i].set_layout, false)))
|
||||
{
|
||||
ERR("Failed to create descriptor set layout %u, hr %#x.", i, hr);
|
||||
goto fail;
|
||||
}
|
||||
|
||||
if (FAILED(hr = vkd3d_create_pipeline_layout(device, 1, set_layouts[i].set_layout,
|
||||
1, &push_constant_range, set_layouts[i].pipeline_layout)))
|
||||
{
|
||||
ERR("Failed to create pipeline layout %u, hr %#x.", i, hr);
|
||||
goto fail;
|
||||
}
|
||||
}
|
||||
|
||||
shader_interface.type = VKD3D_SHADER_STRUCTURE_TYPE_SHADER_INTERFACE_INFO;
|
||||
shader_interface.next = NULL;
|
||||
shader_interface.flags = 0;
|
||||
shader_interface.descriptor_tables.offset = 0;
|
||||
shader_interface.descriptor_tables.count = 0;
|
||||
shader_interface.bindings = &binding;
|
||||
shader_interface.binding_count = 1;
|
||||
shader_interface.push_constant_buffers = &push_constant;
|
||||
shader_interface.push_constant_buffer_count = 1;
|
||||
shader_interface.push_constant_ubo_binding = NULL;
|
||||
|
||||
for (i = 0; i < ARRAY_SIZE(pipelines); ++i)
|
||||
{
|
||||
if (pipelines[i].pipeline_layout == &state->vk_pipeline_layout_buffer)
|
||||
binding.flags = VKD3D_SHADER_BINDING_FLAG_BUFFER;
|
||||
else
|
||||
binding.flags = VKD3D_SHADER_BINDING_FLAG_IMAGE;
|
||||
|
||||
if (FAILED(hr = vkd3d_create_compute_pipeline(device, &pipelines[i].code, &shader_interface,
|
||||
*pipelines[i].pipeline_layout, pipelines[i].pipeline)))
|
||||
{
|
||||
ERR("Failed to create compute pipeline %u, hr %#x.", i, hr);
|
||||
goto fail;
|
||||
}
|
||||
}
|
||||
|
||||
return S_OK;
|
||||
|
||||
fail:
|
||||
vkd3d_uav_clear_state_cleanup(state, device);
|
||||
return hr;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1302,24 +1302,25 @@ struct vkd3d_format_compatibility_list
|
|||
VkFormat vk_formats[VKD3D_MAX_COMPATIBLE_FORMAT_COUNT];
|
||||
};
|
||||
|
||||
struct vkd3d_uav_clear_args
|
||||
/* meta operations */
|
||||
struct vkd3d_clear_uav_args
|
||||
{
|
||||
VkClearColorValue colour;
|
||||
VkClearColorValue clear_color;
|
||||
VkOffset2D offset;
|
||||
VkExtent2D extent;
|
||||
};
|
||||
|
||||
struct vkd3d_uav_clear_pipelines
|
||||
struct vkd3d_clear_uav_pipelines
|
||||
{
|
||||
VkPipeline buffer;
|
||||
VkPipeline image_1d;
|
||||
VkPipeline image_1d_array;
|
||||
VkPipeline image_2d;
|
||||
VkPipeline image_2d_array;
|
||||
VkPipeline image_3d;
|
||||
VkPipeline image_1d_array;
|
||||
VkPipeline image_2d_array;
|
||||
};
|
||||
|
||||
struct vkd3d_uav_clear_state
|
||||
struct vkd3d_clear_uav_ops
|
||||
{
|
||||
VkDescriptorSetLayout vk_set_layout_buffer;
|
||||
VkDescriptorSetLayout vk_set_layout_image;
|
||||
|
|
@ -1327,12 +1328,40 @@ struct vkd3d_uav_clear_state
|
|||
VkPipelineLayout vk_pipeline_layout_buffer;
|
||||
VkPipelineLayout vk_pipeline_layout_image;
|
||||
|
||||
struct vkd3d_uav_clear_pipelines pipelines_float;
|
||||
struct vkd3d_uav_clear_pipelines pipelines_uint;
|
||||
struct vkd3d_clear_uav_pipelines clear_float;
|
||||
struct vkd3d_clear_uav_pipelines clear_uint;
|
||||
};
|
||||
|
||||
HRESULT vkd3d_uav_clear_state_init(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
|
||||
void vkd3d_uav_clear_state_cleanup(struct vkd3d_uav_clear_state *state, struct d3d12_device *device) DECLSPEC_HIDDEN;
|
||||
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_cleanup(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_cleanup(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device) DECLSPEC_HIDDEN;
|
||||
|
||||
struct vkd3d_physical_device_info
|
||||
{
|
||||
|
|
@ -1430,7 +1459,7 @@ struct d3d12_device
|
|||
const struct vkd3d_format_compatibility_list *format_compatibility_lists;
|
||||
struct vkd3d_null_resources null_resources;
|
||||
struct vkd3d_bindless_state bindless_state;
|
||||
struct vkd3d_uav_clear_state uav_clear_state;
|
||||
struct vkd3d_meta_ops meta_ops;
|
||||
};
|
||||
|
||||
HRESULT d3d12_device_create(struct vkd3d_instance *instance,
|
||||
|
|
|
|||
File diff suppressed because it is too large
Load Diff
Loading…
Reference in New Issue