mirrors
/
vkd3d-proton
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

vkd3d: Add meta pipeline to clear raw storage buffers. 

Signed-off-by: Philip Rebohle <philip.rebohle@tu-dortmund.de>
This commit is contained in:
Philip Rebohle 2020-10-12 16:28:35 +02:00 committed by Hans-Kristian Arntzen
parent 290f44254a
commit 3e1445eacb
6 changed files with 41 additions and 7 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

2
libs/vkd3d/command.c
View File

@ -5888,7 +5888,7 @@ static void d3d12_command_list_clear_uav(struct d3d12_command_list *list,
miplevel_idx = 0;
layer_count = 1;
pipeline = vkd3d_meta_get_clear_buffer_uav_pipeline(
&list->device->meta_ops, view->format->type == VKD3D_FORMAT_TYPE_UINT);
&list->device->meta_ops, view->format->type == VKD3D_FORMAT_TYPE_UINT, false);
workgroup_size = vkd3d_meta_get_clear_buffer_uav_workgroup_size();
}

1
libs/vkd3d/meson.build
View File

@ -1,5 +1,6 @@
vkd3d_shaders =[
'shaders/cs_clear_uav_buffer_float.comp',
'shaders/cs_clear_uav_buffer_raw.comp',
'shaders/cs_clear_uav_buffer_uint.comp',
'shaders/cs_clear_uav_image_1d_array_float.comp',
'shaders/cs_clear_uav_image_1d_array_uint.comp',

17
libs/vkd3d/meta.c
View File

@ -334,8 +334,9 @@ HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
}
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 },
{ &meta_clear_uav_ops->vk_set_layout_buffer_raw, &meta_clear_uav_ops->vk_pipeline_layout_buffer_raw, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER },
{ &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 {
@ -367,6 +368,9 @@ HRESULT vkd3d_clear_uav_ops_init(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
{ &meta_clear_uav_ops->clear_uint.buffer,
&meta_clear_uav_ops->vk_pipeline_layout_buffer,
SPIRV_CODE(cs_clear_uav_buffer_uint) },
{ &meta_clear_uav_ops->clear_uint.buffer_raw,
&meta_clear_uav_ops->vk_pipeline_layout_buffer_raw,
SPIRV_CODE(cs_clear_uav_buffer_raw) },
{ &meta_clear_uav_ops->clear_uint.image_1d,
&meta_clear_uav_ops->vk_pipeline_layout_image,
SPIRV_CODE(cs_clear_uav_image_1d_uint) },
@ -444,15 +448,18 @@ void vkd3d_clear_uav_ops_cleanup(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
&meta_clear_uav_ops->clear_uint,
};
VK_CALL(vkDestroyDescriptorSetLayout(device->vk_device, meta_clear_uav_ops->vk_set_layout_buffer_raw, NULL));
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_raw, 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]->buffer_raw, 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));
@ -462,18 +469,18 @@ void vkd3d_clear_uav_ops_cleanup(struct vkd3d_clear_uav_ops *meta_clear_uav_ops,
}
struct vkd3d_clear_uav_pipeline vkd3d_meta_get_clear_buffer_uav_pipeline(struct vkd3d_meta_ops *meta_ops,
bool as_uint)
bool as_uint, bool raw)
{
struct vkd3d_clear_uav_ops *meta_clear_uav_ops = &meta_ops->clear_uav;
struct vkd3d_clear_uav_pipeline info;
const struct vkd3d_clear_uav_pipelines *pipelines = as_uint
const struct vkd3d_clear_uav_pipelines *pipelines = (as_uint || raw)
? &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;
info.vk_pipeline = raw ? pipelines->buffer_raw : pipelines->buffer;
return info;
}

22
libs/vkd3d/shaders/cs_clear_uav_buffer_raw.comp Normal file
View File

@ -0,0 +1,22 @@
#version 450
layout(local_size_x = 128) in;
layout(binding = 0)
writeonly buffer dst_buf {
uint data[];
} 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)
dst.data[u_info.dst_offset.x + thread_id] = u_info.clear_value.x;
}

5
libs/vkd3d/vkd3d_private.h
View File

@ -1644,6 +1644,7 @@ struct vkd3d_clear_uav_args
struct vkd3d_clear_uav_pipelines
{
VkPipeline buffer;
VkPipeline buffer_raw;
VkPipeline image_1d;
VkPipeline image_2d;
VkPipeline image_3d;
@ -1653,9 +1654,11 @@ struct vkd3d_clear_uav_pipelines
struct vkd3d_clear_uav_ops
{
VkDescriptorSetLayout vk_set_layout_buffer_raw;
VkDescriptorSetLayout vk_set_layout_buffer;
VkDescriptorSetLayout vk_set_layout_image;
VkPipelineLayout vk_pipeline_layout_buffer_raw;
VkPipelineLayout vk_pipeline_layout_buffer;
VkPipelineLayout vk_pipeline_layout_image;
@ -1783,7 +1786,7 @@ HRESULT vkd3d_meta_ops_init(struct vkd3d_meta_ops *meta_ops, struct d3d12_device
HRESULT vkd3d_meta_ops_cleanup(struct vkd3d_meta_ops *meta_ops, struct d3d12_device *device);
struct vkd3d_clear_uav_pipeline vkd3d_meta_get_clear_buffer_uav_pipeline(struct vkd3d_meta_ops *meta_ops,
bool as_uint);
bool as_uint, bool raw);
struct vkd3d_clear_uav_pipeline vkd3d_meta_get_clear_image_uav_pipeline(struct vkd3d_meta_ops *meta_ops,
VkImageViewType image_view_type, bool as_uint);
VkExtent3D vkd3d_meta_get_clear_image_uav_workgroup_size(VkImageViewType view_type);

1
libs/vkd3d/vkd3d_shaders.h
View File

@ -30,6 +30,7 @@ enum vkd3d_meta_copy_mode
#include <cs_clear_uav_buffer_float.h>
#include <cs_clear_uav_buffer_uint.h>
#include <cs_clear_uav_buffer_raw.h>
#include <cs_clear_uav_image_1d_array_float.h>
#include <cs_clear_uav_image_1d_array_uint.h>
#include <cs_clear_uav_image_1d_float.h>