From de439c4098b4757ebcd6bd85f3d68723d9f3cfde Mon Sep 17 00:00:00 2001 From: Jorge Pineda Date: Tue, 2 Jul 2024 09:34:06 -0700 Subject: [PATCH] [ET-VK] Delete unused functions in `Context.h`, `Command.h` and `Utils.h` TSIA Differential Revision: [D59281542](https://our.internmc.facebook.com/intern/diff/D59281542/) [ghstack-poisoned] --- backends/vulkan/runtime/api/Command.cpp | 144 ------------------------ backends/vulkan/runtime/api/Command.h | 31 +---- backends/vulkan/runtime/api/Context.h | 127 +-------------------- backends/vulkan/runtime/api/Utils.h | 13 --- 4 files changed, 3 insertions(+), 312 deletions(-) diff --git a/backends/vulkan/runtime/api/Command.cpp b/backends/vulkan/runtime/api/Command.cpp index 9c70cfa60b2..47b640417f5 100644 --- a/backends/vulkan/runtime/api/Command.cpp +++ b/backends/vulkan/runtime/api/Command.cpp @@ -182,150 +182,6 @@ void CommandBuffer::dispatch(const utils::uvec3& global_workgroup_size) { state_ = CommandBuffer::State::RECORDING; } -void CommandBuffer::copy_buffer_to_buffer( - const VulkanBuffer& source, - const VulkanBuffer& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - VK_CHECK_COND( - state_ == CommandBuffer::State::BARRIERS_INSERTED, - "Vulkan CommandBuffer: called copy_buffer_to_buffer() on a command buffer whose state " - "is not BARRIERS_INSERTED."); - - const VkBufferCopy copy_details{ - src_offset.data[0u], // srcOffset - dst_offset.data[0u], // dstOffset - copy_range.data[0u], // size - }; - - vkCmdCopyBuffer( - handle_, source.handle(), destination.handle(), 1u, ©_details); - - state_ = CommandBuffer::State::RECORDING; -} - -void CommandBuffer::copy_texture_to_texture( - const VulkanImage& source, - const VulkanImage& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - VK_CHECK_COND( - state_ == CommandBuffer::State::BARRIERS_INSERTED, - "Vulkan CommandBuffer: called copy_texture_to_texture() on a command buffer whose state " - "is not BARRIERS_INSERTED."); - - const VkImageSubresourceLayers src_subresource_layers{ - VK_IMAGE_ASPECT_COLOR_BIT, // aspectMask - 0u, // mipLevel - 0u, // baseArrayLayer - 1u, // layerCount - }; - - const VkImageSubresourceLayers dst_subresource_layers{ - VK_IMAGE_ASPECT_COLOR_BIT, // aspectMask - 0u, // mipLevel - 0u, // baseArrayLayer - 1u, // layerCount - }; - - const VkImageCopy copy_details{ - src_subresource_layers, // srcSubresource - create_offset3d(src_offset), // srcOffset - dst_subresource_layers, // dstSubresource - create_offset3d(dst_offset), // dstOffset - create_extent3d(copy_range), // extent - }; - - vkCmdCopyImage( - handle_, - source.handle(), - source.layout(), - destination.handle(), - destination.layout(), - 1u, - ©_details); - - state_ = CommandBuffer::State::RECORDING; -} - -void CommandBuffer::copy_texture_to_buffer( - const VulkanImage& source, - const VulkanBuffer& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - VK_CHECK_COND( - state_ == CommandBuffer::State::BARRIERS_INSERTED, - "Vulkan CommandBuffer: called copy_texture_to_buffer() on a command buffer whose state " - "is not BARRIERS_INSERTED."); - - const VkImageSubresourceLayers src_subresource_layers{ - VK_IMAGE_ASPECT_COLOR_BIT, // aspectMask - 0u, // mipLevel - 0u, // baseArrayLayer - 1u, // layerCount - }; - - const VkBufferImageCopy copy_details{ - dst_offset.data[0u], // bufferOffset - dst_offset.data[1u], // bufferRowLength - dst_offset.data[2u], // bufferImageHeight - src_subresource_layers, // imageSubresource - create_offset3d(src_offset), // imageOffset - create_extent3d(copy_range), // imageExtent - }; - - vkCmdCopyImageToBuffer( - handle_, - source.handle(), - source.layout(), - destination.handle(), - 1u, - ©_details); - - state_ = CommandBuffer::State::RECORDING; -} - -void CommandBuffer::copy_buffer_to_texture( - const VulkanBuffer& source, - const VulkanImage& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - VK_CHECK_COND( - state_ == CommandBuffer::State::BARRIERS_INSERTED, - "Vulkan CommandBuffer: called copy_buffer_to_texture() on a command buffer whose state " - "is not BARRIERS_INSERTED."); - - const VkImageSubresourceLayers dst_subresource_layers{ - VK_IMAGE_ASPECT_COLOR_BIT, // aspectMask - 0u, // mipLevel - 0u, // baseArrayLayer - 1u, // layerCount - }; - - const VkBufferImageCopy copy_details{ - src_offset.data[0u], // bufferOffset - src_offset.data[1u], // bufferRowLength - src_offset.data[2u], // bufferImageHeight - dst_subresource_layers, // imageSubresource - create_offset3d(dst_offset), // imageOffset - create_extent3d(copy_range), // imageExtent - }; - - vkCmdCopyBufferToImage( - handle_, - source.handle(), - destination.handle(), - destination.layout(), - 1u, - ©_details); - - state_ = CommandBuffer::State::RECORDING; -} - void CommandBuffer::write_timestamp(VkQueryPool querypool, const uint32_t idx) const { VK_CHECK_COND( diff --git a/backends/vulkan/runtime/api/Command.h b/backends/vulkan/runtime/api/Command.h index ff009de8fc0..c04339ceddd 100644 --- a/backends/vulkan/runtime/api/Command.h +++ b/backends/vulkan/runtime/api/Command.h @@ -39,8 +39,7 @@ class CommandBuffer final { enum State { INVALID, // Used to indicate the command buffer is moved from NEW, // Set during constructor - RECORDING, // Set during call to begin(), dispatch(), and - // copy_*_to_*() + RECORDING, // Set during call to begin() and dispatch() PIPELINE_BOUND, // Set during call to bind_pipeline() DESCRIPTORS_BOUND, // Set during call to bind_descriptors() BARRIERS_INSERTED, // Set during call to insert_barrier() @@ -93,34 +92,6 @@ class CommandBuffer final { void insert_barrier(PipelineBarrier& pipeline_barrier); void dispatch(const utils::uvec3&); - void copy_buffer_to_buffer( - const VulkanBuffer&, - const VulkanBuffer&, - const utils::uvec3&, - const utils::uvec3&, - const utils::uvec3&); - - void copy_texture_to_texture( - const VulkanImage&, - const VulkanImage&, - const utils::uvec3&, - const utils::uvec3&, - const utils::uvec3&); - - void copy_texture_to_buffer( - const VulkanImage&, - const VulkanBuffer&, - const utils::uvec3&, - const utils::uvec3&, - const utils::uvec3&); - - void copy_buffer_to_texture( - const VulkanBuffer&, - const VulkanImage&, - const utils::uvec3&, - const utils::uvec3&, - const utils::uvec3&); - void write_timestamp(VkQueryPool, const uint32_t) const; void reset_querypool(VkQueryPool, const uint32_t, const uint32_t) const; diff --git a/backends/vulkan/runtime/api/Context.h b/backends/vulkan/runtime/api/Context.h index 8d023f73eff..bdbabf82f48 100644 --- a/backends/vulkan/runtime/api/Context.h +++ b/backends/vulkan/runtime/api/Context.h @@ -36,8 +36,8 @@ struct ContextConfig final { // // Vulkan Context holds onto all relevant Vulkan state as it pertains to our -// use of Vulkan in PyTorch. A Context is associated with one, and only one, -// Adapter as a precursor to multi-GPU support. All Vulkan tensors in PyTorch +// use of Vulkan in PyTorch. A Context is associated with one, and only one, +// Adapter as a precursor to multi-GPU support. All Vulkan tensors in PyTorch // are associated with a Context to make tensor <-> device affinity explicit. // The context is currently a global object, but technically it does not need // to be if we were to make it explicit to the user. @@ -200,16 +200,6 @@ class Context final { const ShaderInfo&, const utils::uvec3&); - template - bool submit_copy( - PipelineBarrier&, - const S&, - const D&, - const utils::uvec3&, - const utils::uvec3&, - const utils::uvec3&, - VkFence fence_handle); - template bool submit_compute_job( const ShaderInfo&, @@ -383,119 +373,6 @@ inline void bind( } // namespace detail -template -inline void record_copy( - CommandBuffer& cmd, - const S& source, - const D& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) = delete; - -template <> -inline void record_copy( - CommandBuffer& cmd, - const VulkanBuffer& source, - const VulkanBuffer& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - cmd.copy_buffer_to_buffer( - source, destination, copy_range, src_offset, dst_offset); -} - -template <> -inline void record_copy( - CommandBuffer& cmd, - const VulkanImage& source, - const VulkanImage& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - cmd.copy_texture_to_texture( - source, destination, copy_range, src_offset, dst_offset); -} - -template <> -inline void record_copy( - CommandBuffer& cmd, - const VulkanImage& source, - const VulkanBuffer& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - cmd.copy_texture_to_buffer( - source, destination, copy_range, src_offset, dst_offset); -} - -template <> -inline void record_copy( - CommandBuffer& cmd, - const VulkanBuffer& source, - const VulkanImage& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset) { - cmd.copy_buffer_to_texture( - source, destination, copy_range, src_offset, dst_offset); -} - -/* - Records a GPU data copy into the current command buffer. If the number of - submit_*_job calls exceeds the configured frequency, or if a fence is - provided, then the command buffer is submitted to the GPU for execution. - Returns a bool indicating whether or not the function call resulted in a GPU - queue submission. - */ -template -inline bool Context::submit_copy( - PipelineBarrier& pipeline_barrier, - const S& source, - const D& destination, - const utils::uvec3& copy_range, - const utils::uvec3& src_offset, - const utils::uvec3& dst_offset, - VkFence fence_handle) { - // If any of the provided arguments does not have memory associated with it, - // then exit early as there is no work to be done. However, if a fence has - // been passed the command buffer is not empty, then the current command - // buffer must still be submitted so that the fence can be signaled. - if (!source || !destination) { - if (fence_handle != VK_NULL_HANDLE && submit_count_ > 0) { - submit_cmd_to_gpu(fence_handle); - return true; - } - return false; - } - - // Serialize recording to the shared command buffer. Do not initialize with a - // mutex just yet, since in some cases it will be externally managed. - std::unique_lock cmd_lock; - // Refer to comments in submit_compute_job for explanation. - if (fence_handle == VK_NULL_HANDLE) { - cmd_lock = std::unique_lock(cmd_mutex_); - } - - set_cmd(); - - std::string label = "cmd_copy"; - report_shader_dispatch_start(label, {0, 0, 0}, {0, 0, 0}); - - cmd_.insert_barrier(pipeline_barrier); - - record_copy(cmd_, source, destination, copy_range, src_offset, dst_offset); - - report_shader_dispatch_end(); - - submit_count_++; - if (fence_handle != VK_NULL_HANDLE || - submit_count_ >= config_.cmd_submit_frequency) { - submit_cmd_to_gpu(fence_handle); - return true; - } - return false; -} - /* Records a compute shader dispatch into the current command buffer. If the number of submit_*_job calls exceeds the configured frequency, or if a fence diff --git a/backends/vulkan/runtime/api/Utils.h b/backends/vulkan/runtime/api/Utils.h index d6488cf5976..b327f3153e5 100644 --- a/backends/vulkan/runtime/api/Utils.h +++ b/backends/vulkan/runtime/api/Utils.h @@ -462,19 +462,6 @@ inline int64_t multiply_integers(Iter begin, Iter end) { } // namespace utils -inline bool operator==(const utils::uvec3& _1, const utils::uvec3& _2) { - return ( - _1.data[0u] == _2.data[0u] && _1.data[1u] == _2.data[1u] && - _1.data[2u] == _2.data[2u]); -} - -inline VkOffset3D create_offset3d(const utils::uvec3& offsets) { - return VkOffset3D{ - utils::safe_downcast(offsets.data[0u]), - static_cast(offsets.data[1u]), - static_cast(offsets.data[2u])}; -} - inline VkExtent3D create_extent3d(const utils::uvec3& extents) { return VkExtent3D{extents.data[0u], extents.data[1u], extents.data[2u]}; }