diff --git a/backends/vulkan/runtime/api/Context.cpp b/backends/vulkan/runtime/api/Context.cpp index a51fd4ed526..9a43cf455d6 100644 --- a/backends/vulkan/runtime/api/Context.cpp +++ b/backends/vulkan/runtime/api/Context.cpp @@ -235,5 +235,11 @@ UniformParamsBuffer& UniformParamsBuffer::operator=( return *this; } +ParamsBindList::ParamsBindList( + std::initializer_list init_list) { + bind_infos.resize(init_list.size()); + std::copy(init_list.begin(), init_list.end(), bind_infos.begin()); +} + } // namespace api } // namespace vkcompute diff --git a/backends/vulkan/runtime/api/Context.h b/backends/vulkan/runtime/api/Context.h index 5ecde33596f..d79344dce8d 100644 --- a/backends/vulkan/runtime/api/Context.h +++ b/backends/vulkan/runtime/api/Context.h @@ -244,7 +244,7 @@ class UniformParamsBuffer final { } } - VulkanBuffer& buffer() { + const VulkanBuffer& buffer() const { return vulkan_buffer_; } @@ -264,6 +264,12 @@ class UniformParamsBuffer final { } }; +struct ParamsBindList final { + std::vector bind_infos; + + ParamsBindList(std::initializer_list init_list); +}; + class StorageBuffer final { private: Context* context_p_; @@ -331,6 +337,10 @@ inline void arg_is_empty(bool& any_is_empty, const VulkanImage& image) { any_is_empty = any_is_empty || !image; } +inline void arg_is_empty(bool& any_is_empty, const BufferBindInfo& bind_info) { + any_is_empty = any_is_empty || (bind_info.handle == VK_NULL_HANDLE); +} + /* Reports if any VulkanBuffer or VulkanImage argument in a variadic argument list does not have any memory associated with it. diff --git a/backends/vulkan/runtime/api/Descriptor.cpp b/backends/vulkan/runtime/api/Descriptor.cpp index 25cbaeaa10d..572cc674981 100644 --- a/backends/vulkan/runtime/api/Descriptor.cpp +++ b/backends/vulkan/runtime/api/Descriptor.cpp @@ -15,6 +15,18 @@ namespace vkcompute { namespace api { +// +// BufferBinding +// + +BufferBindInfo::BufferBindInfo() + : handle(VK_NULL_HANDLE), offset(0u), range(0u) {} + +BufferBindInfo::BufferBindInfo(const VulkanBuffer& buffer_p) + : handle(buffer_p.handle()), + offset(buffer_p.mem_offset()), + range(buffer_p.mem_range()) {} + // // DescriptorSet // @@ -66,6 +78,21 @@ DescriptorSet& DescriptorSet::bind( return *this; } +DescriptorSet& DescriptorSet::bind( + const uint32_t idx, + const BufferBindInfo& bind_info) { + DescriptorSet::ResourceBinding binder{}; + binder.binding_idx = idx; // binding_idx + binder.descriptor_type = shader_layout_signature_[idx]; // descriptor_type + binder.is_image = false; // is_image + binder.resource_info.buffer_info.buffer = bind_info.handle; // buffer + binder.resource_info.buffer_info.offset = bind_info.offset; // offset + binder.resource_info.buffer_info.range = bind_info.range; // range + add_binding(binder); + + return *this; +} + DescriptorSet& DescriptorSet::bind( const uint32_t idx, const VulkanImage& image) { diff --git a/backends/vulkan/runtime/api/Descriptor.h b/backends/vulkan/runtime/api/Descriptor.h index 9b9dcda208e..0b6b1cd885a 100644 --- a/backends/vulkan/runtime/api/Descriptor.h +++ b/backends/vulkan/runtime/api/Descriptor.h @@ -20,6 +20,20 @@ namespace vkcompute { namespace api { +/* + * Stores the binding information of a Vulkan Buffer so that the buffer can be + * bound at a later time. This struct should only be used if the buffer to be + * bound is guaranteed to be active at the time of binding. + */ +struct BufferBindInfo final { + VkBuffer handle; + VkDeviceSize offset; + VkDeviceSize range; + + BufferBindInfo(); + BufferBindInfo(const VulkanBuffer& buffer_p); +}; + class DescriptorSet final { public: explicit DescriptorSet(VkDevice, VkDescriptorSet, ShaderLayout::Signature); @@ -50,6 +64,7 @@ class DescriptorSet final { std::vector bindings_; public: + DescriptorSet& bind(const uint32_t, const BufferBindInfo&); DescriptorSet& bind(const uint32_t, const VulkanBuffer&); DescriptorSet& bind(const uint32_t, const VulkanImage&); diff --git a/backends/vulkan/runtime/api/Tensor.cpp b/backends/vulkan/runtime/api/Tensor.cpp index ebd78eac20c..a7055c7f147 100644 --- a/backends/vulkan/runtime/api/Tensor.cpp +++ b/backends/vulkan/runtime/api/Tensor.cpp @@ -140,9 +140,9 @@ vTensor::vTensor( sizes_(sizes.begin(), sizes.end()), gpu_sizes_{calc_gpu_sizes(sizes, memory_layout_, storage_type)}, // Utility Uniform Buffers that can be passed to shaders as arguments - cpu_sizes_uniform_(nullptr), - gpu_sizes_uniform_(nullptr), - extents_uniform_(nullptr), + cpu_sizes_uniform_(), + gpu_sizes_uniform_(), + extents_uniform_(), // Construct Tensor storage storage_( context, @@ -189,33 +189,33 @@ api::VulkanBuffer& vTensor::buffer( return storage_.buffer_; } -std::shared_ptr vTensor::cpu_sizes_ubo() { - if (!cpu_sizes_uniform_) { - cpu_sizes_uniform_.reset(new api::UniformParamsBuffer( - storage_.context_, api::utils::make_whcn_ivec4(sizes_))); +const api::BufferBindInfo vTensor::cpu_sizes_ubo() { + if (!cpu_sizes_uniform_.buffer()) { + cpu_sizes_uniform_ = api::UniformParamsBuffer( + storage_.context_, api::utils::make_whcn_ivec4(sizes_)); } - return cpu_sizes_uniform_; + return api::BufferBindInfo(cpu_sizes_uniform_.buffer()); } -std::shared_ptr vTensor::gpu_sizes_ubo() { - if (!gpu_sizes_uniform_) { - gpu_sizes_uniform_.reset(new api::UniformParamsBuffer( - storage_.context_, api::utils::make_whcn_ivec4(gpu_sizes_))); +const api::BufferBindInfo vTensor::gpu_sizes_ubo() { + if (!gpu_sizes_uniform_.buffer()) { + gpu_sizes_uniform_ = api::UniformParamsBuffer( + storage_.context_, api::utils::make_whcn_ivec4(gpu_sizes_)); } - return gpu_sizes_uniform_; + return api::BufferBindInfo(gpu_sizes_uniform_.buffer()); } -std::shared_ptr vTensor::extents_ubo() { - if (!extents_uniform_) { - extents_uniform_.reset(new api::UniformParamsBuffer( +const api::BufferBindInfo vTensor::extents_ubo() { + if (!extents_uniform_.buffer()) { + extents_uniform_ = api::UniformParamsBuffer( storage_.context_, api::utils::uvec4( {storage_.extents_.data[0], storage_.extents_.data[1], storage_.extents_.data[2], - 1u}))); + 1u})); } - return extents_uniform_; + return api::BufferBindInfo(extents_uniform_.buffer()); } VmaAllocationCreateInfo vTensor::get_allocation_create_info() const { @@ -258,16 +258,16 @@ void vTensor::update_size_metadata(const std::vector& new_sizes) { api::utils::uvec3 virtual_extents = create_image_extents(gpu_sizes_, storage_type(), memory_layout_); - if (cpu_sizes_uniform_) { - cpu_sizes_uniform_->update(api::utils::make_whcn_ivec4(sizes_)); + if (cpu_sizes_uniform_.buffer()) { + cpu_sizes_uniform_.update(api::utils::make_whcn_ivec4(sizes_)); } - if (gpu_sizes_uniform_) { - gpu_sizes_uniform_->update(api::utils::make_whcn_ivec4(gpu_sizes_)); + if (gpu_sizes_uniform_.buffer()) { + gpu_sizes_uniform_.update(api::utils::make_whcn_ivec4(gpu_sizes_)); } - if (extents_uniform_) { - extents_uniform_->update(api::utils::uvec4( + if (extents_uniform_.buffer()) { + extents_uniform_.update(api::utils::uvec4( {virtual_extents.data[0], virtual_extents.data[1], virtual_extents.data[2], diff --git a/backends/vulkan/runtime/api/Tensor.h b/backends/vulkan/runtime/api/Tensor.h index ba9c99c4bf0..3718b6e97d9 100644 --- a/backends/vulkan/runtime/api/Tensor.h +++ b/backends/vulkan/runtime/api/Tensor.h @@ -118,17 +118,17 @@ class vTensor final { // A Vulkan uniform buffer containing the tensor sizes in WHCN that can be // passed into a shader. - std::shared_ptr cpu_sizes_uniform_; + api::UniformParamsBuffer cpu_sizes_uniform_; // A Vulkan uniform buffer containing the GPU tensor sizes in WHCN that can // be passed into a shader. GPU sizes refers to the sizes of the tensor after // padding has been applied to one dimension to align it to the next multiple // of 4. - std::shared_ptr gpu_sizes_uniform_; + api::UniformParamsBuffer gpu_sizes_uniform_; // A Vulkan uniform buffer containing the image extents of the underlying // image texture that can be passed into a shader. - std::shared_ptr extents_uniform_; + api::UniformParamsBuffer extents_uniform_; vTensorStorage storage_; @@ -207,21 +207,21 @@ class vTensor final { * shader. Note that the UBO will be created the first time this function is * called. */ - std::shared_ptr cpu_sizes_ubo(); + const api::BufferBindInfo cpu_sizes_ubo(); /* * Get a uniform buffer object containing the tensor GPU sizes to use in a * compute shader. Note that the UBO will be created the first time this * function is called. */ - std::shared_ptr gpu_sizes_ubo(); + const api::BufferBindInfo gpu_sizes_ubo(); /* * Get a uniform buffer object containing the image extents to use in a * compute shader. Note that the UBO will be created the first time this * function is called. */ - std::shared_ptr extents_ubo(); + const api::BufferBindInfo extents_ubo(); inline size_t numel() const { return api::utils::multiply_integers(sizes()); diff --git a/backends/vulkan/runtime/graph/ComputeGraph.cpp b/backends/vulkan/runtime/graph/ComputeGraph.cpp index 4f17b08f5ec..0c7941d6f52 100644 --- a/backends/vulkan/runtime/graph/ComputeGraph.cpp +++ b/backends/vulkan/runtime/graph/ComputeGraph.cpp @@ -59,6 +59,7 @@ ComputeGraph::ComputeGraph(GraphConfig config) config_.contextConfig)}, shared_objects_{}, values_{}, + param_ubos_{}, prepack_nodes_{}, execute_nodes_{}, inputs_{}, diff --git a/backends/vulkan/runtime/graph/ComputeGraph.h b/backends/vulkan/runtime/graph/ComputeGraph.h index 2e39ed1bdfc..8a6700b6e6a 100644 --- a/backends/vulkan/runtime/graph/ComputeGraph.h +++ b/backends/vulkan/runtime/graph/ComputeGraph.h @@ -93,6 +93,7 @@ class ComputeGraph final { std::unique_ptr context_; std::vector shared_objects_; std::vector values_; + std::vector param_ubos_; std::vector> prepack_nodes_; std::vector> execute_nodes_; @@ -314,9 +315,9 @@ class ComputeGraph final { ValueRef set_output_tensor(const ValueRef idx, const bool use_staging = true); template - inline std::shared_ptr create_params_buffer( - const Block& data) { - return std::make_shared(context_.get(), data); + const api::BufferBindInfo create_params_buffer(const Block& data) { + param_ubos_.emplace_back(api::UniformParamsBuffer(context_.get(), data)); + return api::BufferBindInfo(param_ubos_.back().buffer()); } /* diff --git a/backends/vulkan/runtime/graph/ops/ExecuteNode.cpp b/backends/vulkan/runtime/graph/ops/ExecuteNode.cpp index 347511b782c..5195ec772d8 100644 --- a/backends/vulkan/runtime/graph/ops/ExecuteNode.cpp +++ b/backends/vulkan/runtime/graph/ops/ExecuteNode.cpp @@ -20,7 +20,7 @@ ExecuteNode::ExecuteNode( const api::utils::uvec3& global_workgroup_size, const api::utils::uvec3& local_workgroup_size, const std::vector& args, - const std::vector>& params, + const api::ParamsBindList& params, const ResizeFunction& resize_fn, const std::vector& resize_args, const api::SpecVarList& spec_vars) @@ -47,6 +47,7 @@ void ExecuteNode::encode(ComputeGraph* graph) { uint32_t idx = 0; idx = bind_values_to_descriptor_set( graph, args_, pipeline_barrier, descriptor_set, idx); + bind_params_to_descriptor_set(params_, descriptor_set, idx); context->register_shader_dispatch( diff --git a/backends/vulkan/runtime/graph/ops/ExecuteNode.h b/backends/vulkan/runtime/graph/ops/ExecuteNode.h index 106badb06e0..378588e11dc 100644 --- a/backends/vulkan/runtime/graph/ops/ExecuteNode.h +++ b/backends/vulkan/runtime/graph/ops/ExecuteNode.h @@ -54,7 +54,7 @@ class ExecuteNode final { const api::utils::uvec3& global_workgroup_size, const api::utils::uvec3& local_workgroup_size, const std::vector& args, - const std::vector>& params, + const api::ParamsBindList& params, const ResizeFunction& resize_fn = nullptr, const std::vector& resize_args = {}, const api::SpecVarList& spec_vars = {}); @@ -74,7 +74,7 @@ class ExecuteNode final { const api::utils::uvec3 global_workgroup_size_; const api::utils::uvec3 local_workgroup_size_; const std::vector args_; - std::vector> params_; + const api::ParamsBindList params_; const ResizeFunction resize_fn_; const std::vector resize_args_; const api::SpecVarList spec_vars_; diff --git a/backends/vulkan/runtime/graph/ops/PrepackNode.cpp b/backends/vulkan/runtime/graph/ops/PrepackNode.cpp index 74c593e2caa..8002cf92973 100644 --- a/backends/vulkan/runtime/graph/ops/PrepackNode.cpp +++ b/backends/vulkan/runtime/graph/ops/PrepackNode.cpp @@ -31,7 +31,7 @@ PrepackNode::PrepackNode( const api::utils::uvec3& local_workgroup_size, const ValueRef tref, const ValueRef packed, - const std::vector>& params) + const api::ParamsBindList& params) : shader_(shader), noop_shader_(get_noop_shader(graph, packed)), global_workgroup_size_(global_workgroup_size), diff --git a/backends/vulkan/runtime/graph/ops/PrepackNode.h b/backends/vulkan/runtime/graph/ops/PrepackNode.h index c309200efe7..793665e76cc 100644 --- a/backends/vulkan/runtime/graph/ops/PrepackNode.h +++ b/backends/vulkan/runtime/graph/ops/PrepackNode.h @@ -33,7 +33,7 @@ class PrepackNode final { const api::utils::uvec3& local_workgroup_size, const ValueRef tref, const ValueRef packed, - const std::vector>& params); + const api::ParamsBindList& params); ~PrepackNode() = default; @@ -46,7 +46,7 @@ class PrepackNode final { const api::utils::uvec3 local_workgroup_size_; const ValueRef tref_; const ValueRef packed_; - std::vector> params_; + const api::ParamsBindList params_; private: api::StorageBuffer create_staging_buffer(ComputeGraph* graph); diff --git a/backends/vulkan/runtime/graph/ops/utils/BindingUtils.cpp b/backends/vulkan/runtime/graph/ops/utils/BindingUtils.cpp index 9fec07faa13..158fad3cbba 100644 --- a/backends/vulkan/runtime/graph/ops/utils/BindingUtils.cpp +++ b/backends/vulkan/runtime/graph/ops/utils/BindingUtils.cpp @@ -55,12 +55,12 @@ uint32_t bind_values_to_descriptor_set( } uint32_t bind_params_to_descriptor_set( - std::vector>& params, + const api::ParamsBindList& params, api::DescriptorSet& descriptor_set, const uint32_t base_idx) { uint32_t idx = base_idx; - for (auto& param : params) { - descriptor_set.bind(idx++, param->buffer()); + for (auto& param : params.bind_infos) { + descriptor_set.bind(idx++, param); } return idx; } diff --git a/backends/vulkan/runtime/graph/ops/utils/BindingUtils.h b/backends/vulkan/runtime/graph/ops/utils/BindingUtils.h index 298ed8d76fd..8b3e579c746 100644 --- a/backends/vulkan/runtime/graph/ops/utils/BindingUtils.h +++ b/backends/vulkan/runtime/graph/ops/utils/BindingUtils.h @@ -35,7 +35,7 @@ uint32_t bind_values_to_descriptor_set( // uint32_t bind_params_to_descriptor_set( - std::vector>& params, + const api::ParamsBindList& params, api::DescriptorSet& descriptor_set, const uint32_t base_idx); diff --git a/backends/vulkan/test/utils/test_utils.cpp b/backends/vulkan/test/utils/test_utils.cpp index cbba7c45925..141cac64ea4 100644 --- a/backends/vulkan/test/utils/test_utils.cpp +++ b/backends/vulkan/test/utils/test_utils.cpp @@ -35,8 +35,8 @@ void record_nchw_to_image_op( api::PipelineStage::COMPUTE, api::MemoryAccessType::WRITE), src_buffer, - v_dst.gpu_sizes_ubo()->buffer(), - v_dst.cpu_sizes_ubo()->buffer()); + v_dst.gpu_sizes_ubo(), + v_dst.cpu_sizes_ubo()); } void record_image_to_nchw_op( @@ -55,8 +55,8 @@ void record_image_to_nchw_op( VK_NULL_HANDLE, v_src.image(pipeline_barrier, api::PipelineStage::COMPUTE), dst_buffer, - v_src.gpu_sizes_ubo()->buffer(), - v_src.cpu_sizes_ubo()->buffer()); + v_src.gpu_sizes_ubo(), + v_src.cpu_sizes_ubo()); } void record_conv2d_prepack_weights_op( @@ -96,7 +96,7 @@ void record_conv2d_prepack_weights_op( api::PipelineStage::COMPUTE, api::MemoryAccessType::WRITE), src_buffer, - v_dst.gpu_sizes_ubo()->buffer(), + v_dst.gpu_sizes_ubo(), original_sizes_ubo.buffer(), padded_sizes_ubo.buffer()); } @@ -125,7 +125,7 @@ void record_binary_op( api::MemoryAccessType::WRITE), v_in1.image(pipeline_barrier, api::PipelineStage::COMPUTE), v_in2.image(pipeline_barrier, api::PipelineStage::COMPUTE), - v_dst.extents_ubo()->buffer()); + v_dst.extents_ubo()); } void execute_and_check_add( diff --git a/backends/vulkan/test/vulkan_compute_api_test.cpp b/backends/vulkan/test/vulkan_compute_api_test.cpp index 0d89d618166..777f9adab08 100644 --- a/backends/vulkan/test/vulkan_compute_api_test.cpp +++ b/backends/vulkan/test/vulkan_compute_api_test.cpp @@ -923,8 +923,8 @@ void run_from_gpu_test( pipeline_barrier, api::PipelineStage::COMPUTE, api::MemoryAccessType::WRITE), - vten.gpu_sizes_ubo()->buffer(), - vten.cpu_sizes_ubo()->buffer()); + vten.gpu_sizes_ubo(), + vten.cpu_sizes_ubo()); } api::StorageBuffer staging_buffer(api::context(), dtype, vten.gpu_numel());