From 740c12733e7a61f6c3d69186939b5b5e7939fc2c Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Mon, 25 May 2026 16:03:44 +0200 Subject: [PATCH 01/11] vulkan: add optimized device to device copy function --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 651 +++++++++++++++++++++++++-- 1 file changed, 618 insertions(+), 33 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index c1d86aaac5..ef5fb4b91d 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -787,6 +787,7 @@ struct vk_device_struct { uint64_t suballocation_block_size; uint64_t min_imported_host_pointer_alignment; bool external_memory_host {}; + bool external_memory_dma_buf {}; bool fp16; bool bf16; bool pipeline_robustness; @@ -1110,35 +1111,7 @@ struct vk_device_struct { std::unique_ptr memory_logger; - ~vk_device_struct() { - VK_LOG_DEBUG("destroy device " << name); - - device.destroyFence(fence); - - ggml_vk_destroy_buffer(sync_staging); - - if (compute_queue) compute_queue->cmd_pool.destroy(device); - if (transfer_queue) transfer_queue->cmd_pool.destroy(device); - - // Explicitly clear to ensure queues drop their shared_ptrs to handles - // before the Vulkan logical device instance is destroyed - compute_queue.reset(); - transfer_queue.reset(); - - for (auto& pipeline : all_pipelines) { - if (pipeline.expired()) { - continue; - } - - vk_pipeline pl = pipeline.lock(); - ggml_vk_destroy_pipeline(device, pl); - } - all_pipelines.clear(); - - device.destroyDescriptorSetLayout(dsl); - - device.destroy(); - } + ~vk_device_struct(); }; void vk_command_pool::init(vk_device& device, vk_queue *q_) { @@ -1229,6 +1202,25 @@ struct vk_buffer_struct { } }; +#ifdef __linux__ +enum vk_d2d_method { + D2D_UNTESTED, + D2D_DMABUF_P2P, + D2D_DMABUF_GTT, + D2D_SHARED_STAGING, + D2D_STAGING, +}; + +struct vk_d2d_path { + vk_d2d_method method = D2D_UNTESTED; + bool reverse_direction = false; + vk_buffer buf_a; + vk_buffer buf_b; + void * host_ptr = nullptr; + size_t size = 0; +}; +#endif + struct vk_subbuffer { vk_buffer buffer; uint64_t offset; @@ -2564,11 +2556,92 @@ struct vk_instance_t { std::vector device_indices; std::vector device_supports_membudget; vk_device devices[GGML_VK_MAX_DEVICES]; + + ~vk_instance_t(); }; +#ifdef __linux__ +static std::mutex vk_d2d_cache_mutex; +static std::map, vk_d2d_path> vk_d2d_cache; +#endif + static bool vk_instance_initialized = false; static vk_instance_t vk_instance; +vk_instance_t::~vk_instance_t() { +#ifdef __linux__ + { + std::lock_guard guard(vk_d2d_cache_mutex); + for (auto& entry : vk_d2d_cache) { + if (entry.second.host_ptr) { + free(entry.second.host_ptr); + entry.second.host_ptr = nullptr; + } + if (entry.second.buf_a) { + entry.second.buf_a->size = 0; + } + if (entry.second.buf_b) { + entry.second.buf_b->size = 0; + } + } + vk_d2d_cache.clear(); + } +#endif +} + +vk_device_struct::~vk_device_struct() { + VK_LOG_DEBUG("destroy device " << name); + +#ifdef __linux__ + { + std::lock_guard guard(vk_d2d_cache_mutex); + for (auto it = vk_d2d_cache.begin(); it != vk_d2d_cache.end(); ) { + if (it->first.first == this || it->first.second == this) { + if (it->second.host_ptr) { + free(it->second.host_ptr); + it->second.host_ptr = nullptr; + } + if (it->second.buf_a) { + it->second.buf_a->size = 0; + } + if (it->second.buf_b) { + it->second.buf_b->size = 0; + } + it = vk_d2d_cache.erase(it); + } else { + ++it; + } + } + } +#endif + + device.destroyFence(fence); + + ggml_vk_destroy_buffer(sync_staging); + + if (compute_queue) compute_queue->cmd_pool.destroy(device); + if (transfer_queue) transfer_queue->cmd_pool.destroy(device); + + // Explicitly clear to ensure queues drop their shared_ptrs to handles + // before the Vulkan logical device instance is destroyed + compute_queue.reset(); + transfer_queue.reset(); + + for (auto& pipeline : all_pipelines) { + if (pipeline.expired()) { + continue; + } + + vk_pipeline pl = pipeline.lock(); + ggml_vk_destroy_pipeline(device, pl); + } + all_pipelines.clear(); + + device.destroyDescriptorSetLayout(dsl); + + device.destroy(); +} + #ifdef GGML_VULKAN_CHECK_RESULTS static size_t vk_skip_checks; static size_t vk_output_tensor; @@ -3626,6 +3699,258 @@ static void ggml_vk_destroy_buffer(vk_buffer& buf) { buf.reset(); } +#ifdef __linux__ +#include + +static vk_buffer ggml_vk_create_buffer_dma_buf_export(vk_device& device, size_t size, bool device_local) { + vk_buffer buf = std::make_shared(); + + vk::BufferUsageFlags usage_flags = vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eTransferDst; + + vk::ExternalMemoryBufferCreateInfo external_memory_bci; + external_memory_bci.handleTypes = vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT; + + vk::BufferCreateInfo buffer_create_info{ + vk::BufferCreateFlags(), + size, + usage_flags, + vk::SharingMode::eExclusive, + 0, + nullptr, + }; + buffer_create_info.setPNext(&external_memory_bci); + + try { + buf->buffer = device->device.createBuffer(buffer_create_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_create_buffer_dma_buf_export: createBuffer failed: " << e.what()); + return {}; + } + + vk::MemoryRequirements mem_req = device->device.getBufferMemoryRequirements(buf->buffer); + vk::PhysicalDeviceMemoryProperties mem_props = device->physical_device.getMemoryProperties(); + + vk::MemoryPropertyFlags req_flags; + if (device_local) { + req_flags = vk::MemoryPropertyFlagBits::eDeviceLocal; + } else { + req_flags = vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent; + } + + const std::vector memory_type_indices = ggml_vk_find_memory_properties(&mem_props, &mem_req, req_flags); + if (memory_type_indices.empty()) { + VK_LOG_DEBUG("ggml_vk_create_buffer_dma_buf_export: no suitable memory type"); + device->device.destroyBuffer(buf->buffer); + return {}; + } + + // For GTT, prefer non-device-local host-visible memory + uint32_t chosen_idx = memory_type_indices[0]; + if (!device_local) { + for (uint32_t idx : memory_type_indices) { + if (!(mem_props.memoryTypes[idx].propertyFlags & vk::MemoryPropertyFlagBits::eDeviceLocal)) { + chosen_idx = idx; + break; + } + } + } + + vk::ExportMemoryAllocateInfo export_info; + export_info.handleTypes = vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT; + + vk::MemoryAllocateInfo alloc_info; + alloc_info.allocationSize = mem_req.size; + alloc_info.memoryTypeIndex = chosen_idx; + alloc_info.setPNext(&export_info); + + try { + buf->device_memory = device->device.allocateMemory(alloc_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_create_buffer_dma_buf_export: allocateMemory failed: " << e.what()); + device->device.destroyBuffer(buf->buffer); + return {}; + } + + buf->memory_property_flags = mem_props.memoryTypes[chosen_idx].propertyFlags; + buf->ptr = nullptr; + + device->device.bindBufferMemory(buf->buffer, buf->device_memory, 0); + buf->device = device; + buf->size = size; + + device->memory_logger->log_allocation(buf, size); + + return buf; +} + +static int ggml_vk_export_dma_buf_fd(vk_device& device, vk_buffer& buf) { + vk::MemoryGetFdInfoKHR fd_info; + fd_info.memory = buf->device_memory; + fd_info.handleType = vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT; + + try { + return device->device.getMemoryFdKHR(fd_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_export_dma_buf_fd: getMemoryFdKHR failed: " << e.what()); + return -1; + } +} + +static vk_buffer ggml_vk_import_dma_buf_fd(vk_device& device, int fd, size_t size, vk::MemoryPropertyFlags req_flags) { + vk_buffer buf = std::make_shared(); + + vk::ExternalMemoryBufferCreateInfo external_memory_bci; + external_memory_bci.handleTypes = vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT; + + vk::BufferCreateInfo buffer_create_info{ + vk::BufferCreateFlags(), + size, + vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eTransferDst, + vk::SharingMode::eExclusive, + 0, + nullptr, + }; + buffer_create_info.setPNext(&external_memory_bci); + + try { + buf->buffer = device->device.createBuffer(buffer_create_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_import_dma_buf_fd: createBuffer failed: " << e.what()); + close(fd); + return {}; + } + + vk::MemoryRequirements mem_req = device->device.getBufferMemoryRequirements(buf->buffer); + vk::PhysicalDeviceMemoryProperties mem_props = device->physical_device.getMemoryProperties(); + + vk::MemoryFdPropertiesKHR fd_props; + try { + fd_props = device->device.getMemoryFdPropertiesKHR(vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT, fd); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_import_dma_buf_fd: getMemoryFdPropertiesKHR failed: " << e.what()); + device->device.destroyBuffer(buf->buffer); + close(fd); + return {}; + } + + uint32_t memory_type_idx; + for (memory_type_idx = 0; memory_type_idx < mem_props.memoryTypeCount; ++memory_type_idx) { + if (!(fd_props.memoryTypeBits & (1u << memory_type_idx))) { + continue; + } + if (!(mem_req.memoryTypeBits & (1u << memory_type_idx))) { + continue; + } + if ((mem_props.memoryTypes[memory_type_idx].propertyFlags & req_flags) == req_flags) { + break; + } + } + if (memory_type_idx == mem_props.memoryTypeCount) { + VK_LOG_DEBUG("ggml_vk_import_dma_buf_fd: no suitable memory type"); + device->device.destroyBuffer(buf->buffer); + close(fd); + return {}; + } + + vk::ImportMemoryFdInfoKHR import_info; + import_info.handleType = vk::ExternalMemoryHandleTypeFlagBits::eDmaBufEXT; + import_info.fd = fd; + + vk::MemoryAllocateInfo alloc_info; + alloc_info.allocationSize = mem_req.size; + alloc_info.memoryTypeIndex = memory_type_idx; + alloc_info.setPNext(&import_info); + + try { + buf->device_memory = device->device.allocateMemory(alloc_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_import_dma_buf_fd: allocateMemory failed: " << e.what()); + device->device.destroyBuffer(buf->buffer); + close(fd); + return {}; + } + // fd is consumed by successful import — do not close + + buf->memory_property_flags = mem_props.memoryTypes[memory_type_idx].propertyFlags; + buf->ptr = nullptr; + + device->device.bindBufferMemory(buf->buffer, buf->device_memory, 0); + buf->device = device; + buf->size = size; + + device->memory_logger->log_allocation(buf, size); + + return buf; +} + +static bool ggml_vk_d2d_try_dma_buf(vk_device& exporter, vk_device& importer, size_t size, bool device_local, + vk_buffer& out_export_buf, vk_buffer& out_import_buf) { + out_export_buf = ggml_vk_create_buffer_dma_buf_export(exporter, size, device_local); + if (!out_export_buf) { + return false; + } + + int fd = ggml_vk_export_dma_buf_fd(exporter, out_export_buf); + if (fd < 0) { + ggml_vk_destroy_buffer(out_export_buf); + return false; + } + + vk::MemoryPropertyFlags import_flags = device_local + ? vk::MemoryPropertyFlagBits::eDeviceLocal + : (vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent); + + out_import_buf = ggml_vk_import_dma_buf_fd(importer, fd, size, import_flags); + if (!out_import_buf) { + ggml_vk_destroy_buffer(out_export_buf); + return false; + } + + return true; +} + +static bool ggml_vk_d2d_try_shared_staging(vk_device& dev_a, vk_device& dev_b, size_t size, + vk_buffer& out_buf_a, vk_buffer& out_buf_b, void*& out_host_ptr) { + if (!dev_a->external_memory_host || !dev_b->external_memory_host) { + return false; + } + + uint64_t align = std::max(dev_a->min_imported_host_pointer_alignment, dev_b->min_imported_host_pointer_alignment); + size_t alloc_size = (size + align - 1) & ~(align - 1); + + void * ptr = nullptr; + if (posix_memalign(&ptr, align, alloc_size) != 0 || ptr == nullptr) { + return false; + } + + const vk::MemoryPropertyFlags flags = vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent | vk::MemoryPropertyFlagBits::eHostCached; + + try { + out_buf_a = ggml_vk_create_buffer(dev_a, alloc_size, {flags}, ptr); + } catch (const vk::SystemError&) { + out_buf_a = {}; + } + if (!out_buf_a) { + free(ptr); + return false; + } + + try { + out_buf_b = ggml_vk_create_buffer(dev_b, alloc_size, {flags}, ptr); + } catch (const vk::SystemError&) { + out_buf_b = {}; + } + if (!out_buf_b) { + ggml_vk_destroy_buffer(out_buf_a); + free(ptr); + return false; + } + + out_host_ptr = ptr; + return true; +} +#endif + static vk_subbuffer ggml_vk_subbuffer(const ggml_backend_vk_context* ctx, const vk_buffer& buf, size_t offset = 0) { return { buf, offset, ggml_vk_get_max_buffer_range(ctx, buf, offset) }; } @@ -6211,6 +6536,10 @@ static vk_device ggml_vk_get_device(size_t idx) { bool dot2_f16_support = false; bool ocp_microscaling_extension = false; bool shader_float8_extension = false; +#ifdef __linux__ + bool dma_buf_support = false; + bool external_memory_fd_support = false; +#endif for (const auto& properties : ext_props) { if (strcmp("VK_KHR_maintenance4", properties.extensionName) == 0) { @@ -6271,6 +6600,12 @@ static vk_device ggml_vk_get_device(size_t idx) { device->memory_priority = true; } else if (strcmp("VK_EXT_external_memory_host", properties.extensionName) == 0) { device->external_memory_host = true; +#ifdef __linux__ + } else if (strcmp("VK_EXT_external_memory_dma_buf", properties.extensionName) == 0) { + dma_buf_support = true; + } else if (strcmp("VK_KHR_external_memory_fd", properties.extensionName) == 0) { + external_memory_fd_support = true; +#endif #if defined(VK_EXT_shader_64bit_indexing) } else if (strcmp("VK_EXT_shader_64bit_indexing", properties.extensionName) == 0) { device->shader_64b_indexing = true; @@ -6450,6 +6785,10 @@ static vk_device ggml_vk_get_device(size_t idx) { device->min_imported_host_pointer_alignment = external_memory_host_props.minImportedHostPointerAlignment; +#ifdef __linux__ + device->external_memory_dma_buf = dma_buf_support && external_memory_fd_support; +#endif + device->max_workgroup_size_log2 = uint32_t(log2f(float(device->properties.limits.maxComputeWorkGroupInvocations))); std::vector queue_family_props = device->physical_device.getQueueFamilyProperties(); @@ -6623,6 +6962,13 @@ static vk_device ggml_vk_get_device(size_t idx) { device_extensions.push_back("VK_EXT_external_memory_host"); } +#ifdef __linux__ + if (device->external_memory_dma_buf) { + device_extensions.push_back("VK_EXT_external_memory_dma_buf"); + device_extensions.push_back("VK_KHR_external_memory_fd"); + } +#endif + #if defined(VK_EXT_shader_64bit_indexing) VkPhysicalDeviceShader64BitIndexingFeaturesEXT shader_64bit_indexing_features {}; shader_64bit_indexing_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_SHADER_64_BIT_INDEXING_FEATURES_EXT; @@ -8651,6 +8997,232 @@ static void ggml_vk_buffer_read(vk_buffer& src, size_t offset, void * dst, size_ ggml_vk_buffer_read_2d(src, offset, dst, size, size, size, 1); } +#ifdef __linux__ +static bool ggml_vk_d2d_test_copy(vk_device& device, vk_buffer& shared_buf, size_t size) { + vk_buffer tmp; + try { + tmp = ggml_vk_create_buffer(device, size, + {vk::MemoryPropertyFlagBits::eDeviceLocal, + vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent}); + } catch (...) { + return false; + } + if (!tmp) { + return false; + } + + std::lock_guard guard(device->mutex); + vk_context subctx = ggml_vk_create_temporary_context(device->transfer_queue.cmd_pool); + ggml_vk_ctx_begin(device, subctx); + VkBufferCopy bc{ 0, 0, size }; + vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)shared_buf->buffer, (VkBuffer)tmp->buffer, 1, &bc); + ggml_vk_ctx_end(subctx); + try { + ggml_vk_submit(subctx, device->fence); + VK_CHECK(device->device.waitForFences({ device->fence }, true, UINT64_MAX), "d2d test copy waitForFences"); + device->device.resetFences({ device->fence }); + ggml_vk_queue_command_pools_cleanup(device); + ggml_vk_destroy_buffer(tmp); + return true; + } catch (...) { + try { + device->device.resetFences({ device->fence }); + ggml_vk_queue_command_pools_cleanup(device); + } catch (...) {} + ggml_vk_destroy_buffer(tmp); + return false; + } +} + +static const size_t VK_D2D_PROBE_SIZE = 4096; + +static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev) { + VK_LOG_DEBUG("ggml_vk_probe_d2d_path(" << src_dev->name << " -> " << dst_dev->name << ")"); + vk_d2d_path path; + + bool src_nvidia = src_dev->vendor_id == 0x10de; + bool dst_nvidia = dst_dev->vendor_id == 0x10de; + bool cross_vendor_nvidia = src_nvidia != dst_nvidia; + + // 1. dmabuf_p2p — skip if cross-vendor NVIDIA + if (src_dev->external_memory_dma_buf && dst_dev->external_memory_dma_buf && !cross_vendor_nvidia) { + // Try src exports VRAM, dst imports (read direction) + vk_buffer exp_buf, imp_buf; + if (ggml_vk_d2d_try_dma_buf(src_dev, dst_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_P2P; + path.reverse_direction = false; + path.buf_a = exp_buf; + path.buf_b = imp_buf; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } + + // Try dst exports VRAM, src imports (write direction) + if (ggml_vk_d2d_try_dma_buf(dst_dev, src_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_P2P; + path.reverse_direction = true; + path.buf_a = exp_buf; + path.buf_b = imp_buf; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } + } + + // 2. dmabuf_gtt + if (src_dev->external_memory_dma_buf && dst_dev->external_memory_dma_buf) { + vk_buffer exp_buf, imp_buf; + // Try src exports GTT + if (ggml_vk_d2d_try_dma_buf(src_dev, dst_dev, VK_D2D_PROBE_SIZE, false, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(src_dev, exp_buf, VK_D2D_PROBE_SIZE) && + ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_GTT; + path.reverse_direction = false; + path.buf_a = exp_buf; + path.buf_b = imp_buf; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (src exports GTT)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } + + // Try dst exports GTT + if (ggml_vk_d2d_try_dma_buf(dst_dev, src_dev, VK_D2D_PROBE_SIZE, false, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(dst_dev, exp_buf, VK_D2D_PROBE_SIZE) && + ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_GTT; + path.reverse_direction = true; + path.buf_a = exp_buf; + path.buf_b = imp_buf; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (dst exports GTT)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } + } + + // 3. shared_staging + { + vk_buffer buf_a, buf_b; + void * host_ptr = nullptr; + if (ggml_vk_d2d_try_shared_staging(src_dev, dst_dev, VK_D2D_PROBE_SIZE, buf_a, buf_b, host_ptr)) { + if (ggml_vk_d2d_test_copy(src_dev, buf_a, VK_D2D_PROBE_SIZE) && + ggml_vk_d2d_test_copy(dst_dev, buf_b, VK_D2D_PROBE_SIZE)) { + path.method = D2D_SHARED_STAGING; + path.buf_a = buf_a; + path.buf_b = buf_b; + path.host_ptr = host_ptr; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: shared_staging\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; + } + ggml_vk_destroy_buffer(buf_a); + ggml_vk_destroy_buffer(buf_b); + free(host_ptr); + } + } + + // 4. Fallback + path.method = D2D_STAGING; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: staging (fallback)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return path; +} + +static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, size_t needed) { + VK_LOG_DEBUG("ggml_vk_d2d_grow_path(" << needed << ", current=" << path.size << ")"); + + vk_buffer new_buf_a, new_buf_b; + void * new_host_ptr = nullptr; + bool ok = false; + + switch (path.method) { + case D2D_DMABUF_P2P: { + vk_device& exporter = path.reverse_direction ? dst_dev : src_dev; + vk_device& importer = path.reverse_direction ? src_dev : dst_dev; + ok = ggml_vk_d2d_try_dma_buf(exporter, importer, needed, true, new_buf_a, new_buf_b); + break; + } + case D2D_DMABUF_GTT: { + vk_device& exporter = path.reverse_direction ? dst_dev : src_dev; + vk_device& importer = path.reverse_direction ? src_dev : dst_dev; + ok = ggml_vk_d2d_try_dma_buf(exporter, importer, needed, false, new_buf_a, new_buf_b); + break; + } + case D2D_SHARED_STAGING: + ok = ggml_vk_d2d_try_shared_staging(src_dev, dst_dev, needed, new_buf_a, new_buf_b, new_host_ptr); + break; + default: + return false; + } + + if (!ok) { + return false; + } + + // Destroy old buffers + ggml_vk_destroy_buffer(path.buf_a); + ggml_vk_destroy_buffer(path.buf_b); + if (path.host_ptr) { + free(path.host_ptr); + } + + path.buf_a = new_buf_a; + path.buf_b = new_buf_b; + path.host_ptr = new_host_ptr; + path.size = needed; + return true; +} + +static vk_d2d_path& ggml_vk_get_d2d_path(vk_device& src_dev, vk_device& dst_dev, size_t size) { + std::lock_guard guard(vk_d2d_cache_mutex); + auto key = std::make_pair(src_dev.get(), dst_dev.get()); + auto it = vk_d2d_cache.find(key); + + if (it == vk_d2d_cache.end()) { + vk_d2d_cache[key] = ggml_vk_probe_d2d_path(src_dev, dst_dev); + it = vk_d2d_cache.find(key); + } + + vk_d2d_path& path = it->second; + + if (path.method != D2D_STAGING && path.size < size) { + if (!ggml_vk_d2d_grow_path(path, src_dev, dst_dev, size)) { + GGML_LOG_WARN("ggml_vulkan: d2d grow failed for %s -> %s, falling back to staging\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_destroy_buffer(path.buf_a); + ggml_vk_destroy_buffer(path.buf_b); + if (path.host_ptr) { + free(path.host_ptr); + path.host_ptr = nullptr; + } + path.method = D2D_STAGING; + path.size = 0; + } + } + + return path; +} +#endif + static void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) { VK_LOG_DEBUG("ggml_vk_buffer_copy_async(" << size << ")"); // Make sure both buffers are on same device @@ -8676,12 +9248,25 @@ static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& sr ggml_vk_queue_command_pools_cleanup(src->device); } else { VK_LOG_DEBUG("ggml_vk_buffer_copy(MULTI_DEVICE, " << size << ")"); - // Copy device to device - ggml_vk_ensure_sync_staging_buffer(src->device, size); +#ifdef __linux__ + vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); - // Copy to src staging buffer + if (path.method != D2D_STAGING) { + // buf_a is on the src-side device, buf_b is on the dst-side device + // For reverse_direction (dst exports), buf_a is on dst_dev and buf_b is on src_dev + vk_buffer& src_side_buf = path.reverse_direction ? path.buf_b : path.buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? path.buf_a : path.buf_b; + + // Hop 1: src GPU copies VRAM -> shared buffer (same-device copy on src) + ggml_vk_buffer_copy(src_side_buf, 0, src, src_offset, size); + // Hop 2: dst GPU copies shared buffer -> VRAM (same-device copy on dst) + ggml_vk_buffer_copy(dst, dst_offset, dst_side_buf, 0, size); + return; + } +#endif + // Fallback: staging with CPU memcpy + ggml_vk_ensure_sync_staging_buffer(src->device, size); ggml_vk_buffer_copy(src->device->sync_staging, 0, src, src_offset, size); - // Copy to dst buffer ggml_vk_buffer_write(dst, dst_offset, src->device->sync_staging->ptr, size); } } From bf316ac516abf598290f61f0dbab793ca6093d5f Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Mon, 25 May 2026 16:27:17 +0200 Subject: [PATCH 02/11] add async copy --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 263 ++++++++++++++++++++++++++- 1 file changed, 262 insertions(+), 1 deletion(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index ef5fb4b91d..c8e4dff0b8 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -788,6 +788,7 @@ struct vk_device_struct { uint64_t min_imported_host_pointer_alignment; bool external_memory_host {}; bool external_memory_dma_buf {}; + bool external_semaphore_fd {}; bool fp16; bool bf16; bool pipeline_robustness; @@ -1218,6 +1219,18 @@ struct vk_d2d_path { vk_buffer buf_b; void * host_ptr = nullptr; size_t size = 0; + + bool async_capable = false; + vk::Semaphore sem_src = VK_NULL_HANDLE; + vk::Semaphore sem_dst = VK_NULL_HANDLE; + uint64_t sem_value = 0; + vk_device_struct * sem_src_device = nullptr; + vk_device_struct * sem_dst_device = nullptr; + + vk_command_pool hop1_cmd_pool; + vk::Fence hop1_fence = VK_NULL_HANDLE; + bool hop1_fence_pending = false; + vk_device_struct * hop1_device = nullptr; }; #endif @@ -2568,11 +2581,16 @@ static std::map, vk_d2d_path> vk static bool vk_instance_initialized = false; static vk_instance_t vk_instance; +#ifdef __linux__ +static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path); +#endif + vk_instance_t::~vk_instance_t() { #ifdef __linux__ { std::lock_guard guard(vk_d2d_cache_mutex); for (auto& entry : vk_d2d_cache) { + ggml_vk_d2d_destroy_shared_semaphore(entry.second); if (entry.second.host_ptr) { free(entry.second.host_ptr); entry.second.host_ptr = nullptr; @@ -2597,6 +2615,7 @@ vk_device_struct::~vk_device_struct() { std::lock_guard guard(vk_d2d_cache_mutex); for (auto it = vk_d2d_cache.begin(); it != vk_d2d_cache.end(); ) { if (it->first.first == this || it->first.second == this) { + ggml_vk_d2d_destroy_shared_semaphore(it->second); if (it->second.host_ptr) { free(it->second.host_ptr); it->second.host_ptr = nullptr; @@ -3949,6 +3968,154 @@ static bool ggml_vk_d2d_try_shared_staging(vk_device& dev_a, vk_device& dev_b, s out_host_ptr = ptr; return true; } + +static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path) { + if (!path.async_capable) { + return; + } + + if (path.hop1_fence_pending && path.hop1_device) { + try { + VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), + "d2d destroy wait hop1 fence"); + } catch (...) {} + path.hop1_fence_pending = false; + } + + if (path.hop1_fence) { + path.hop1_device->device.destroyFence(path.hop1_fence); + path.hop1_fence = VK_NULL_HANDLE; + } + if (path.hop1_cmd_pool.pool) { + path.hop1_cmd_pool.destroy(path.hop1_device->device); + } + + if (path.sem_src) { + path.sem_src_device->device.destroySemaphore(path.sem_src); + path.sem_src = VK_NULL_HANDLE; + } + if (path.sem_dst) { + path.sem_dst_device->device.destroySemaphore(path.sem_dst); + path.sem_dst = VK_NULL_HANDLE; + } + + path.async_capable = false; + path.sem_value = 0; + path.sem_src_device = nullptr; + path.sem_dst_device = nullptr; + path.hop1_device = nullptr; +} + +static bool ggml_vk_d2d_check_timeline_semaphore_export(vk_device& dev) { + vk::SemaphoreTypeCreateInfo sem_type; + sem_type.semaphoreType = vk::SemaphoreType::eTimeline; + + vk::PhysicalDeviceExternalSemaphoreInfo ext_sem_info; + ext_sem_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; + ext_sem_info.pNext = &sem_type; + + vk::ExternalSemaphoreProperties ext_sem_props = dev->physical_device.getExternalSemaphoreProperties(ext_sem_info); + + return (ext_sem_props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eExportable) && + (ext_sem_props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eImportable); +} + +static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { + if (!src_dev->external_semaphore_fd || !dst_dev->external_semaphore_fd) { + return false; + } + + if (!ggml_vk_d2d_check_timeline_semaphore_export(src_dev) || + !ggml_vk_d2d_check_timeline_semaphore_export(dst_dev)) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: timeline semaphore export/import not supported"); + return false; + } + + vk::ExportSemaphoreCreateInfo export_ci; + export_ci.handleTypes = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; + + vk::SemaphoreTypeCreateInfo type_ci; + type_ci.semaphoreType = vk::SemaphoreType::eTimeline; + type_ci.initialValue = 0; + type_ci.pNext = &export_ci; + + vk::SemaphoreCreateInfo sem_ci; + sem_ci.pNext = &type_ci; + + vk::Semaphore src_sem; + try { + src_sem = src_dev->device.createSemaphore(sem_ci); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: createSemaphore on src failed: " << e.what()); + return false; + } + + vk::SemaphoreGetFdInfoKHR get_fd_info; + get_fd_info.semaphore = src_sem; + get_fd_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; + + int fd = -1; + try { + fd = src_dev->device.getSemaphoreFdKHR(get_fd_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: getSemaphoreFdKHR failed: " << e.what()); + src_dev->device.destroySemaphore(src_sem); + return false; + } + if (fd < 0) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: getSemaphoreFdKHR returned invalid fd"); + src_dev->device.destroySemaphore(src_sem); + return false; + } + + vk::SemaphoreTypeCreateInfo dst_type_ci{ vk::SemaphoreType::eTimeline, 0 }; + + vk::SemaphoreCreateInfo dst_sem_ci{}; + dst_sem_ci.pNext = &dst_type_ci; + + vk::Semaphore dst_sem; + try { + dst_sem = dst_dev->device.createSemaphore(dst_sem_ci); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: createSemaphore on dst failed: " << e.what()); + close(fd); + src_dev->device.destroySemaphore(src_sem); + return false; + } + + vk::ImportSemaphoreFdInfoKHR import_info; + import_info.semaphore = dst_sem; + import_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; + import_info.fd = fd; + + try { + dst_dev->device.importSemaphoreFdKHR(import_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: importSemaphoreFdKHR failed: " << e.what()); + close(fd); + dst_dev->device.destroySemaphore(dst_sem); + src_dev->device.destroySemaphore(src_sem); + return false; + } + // fd ownership transferred to driver on successful import + + path.sem_src = src_sem; + path.sem_dst = dst_sem; + path.sem_value = 0; + path.sem_src_device = src_dev.get(); + path.sem_dst_device = dst_dev.get(); + + path.hop1_cmd_pool.init(src_dev, &src_dev->transfer_queue); + path.hop1_fence = src_dev->device.createFence({}); + path.hop1_fence_pending = false; + path.hop1_device = src_dev.get(); + + path.async_capable = true; + + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: async semaphore created\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + return true; +} #endif static vk_subbuffer ggml_vk_subbuffer(const ggml_backend_vk_context* ctx, const vk_buffer& buf, size_t offset = 0) { @@ -6539,6 +6706,8 @@ static vk_device ggml_vk_get_device(size_t idx) { #ifdef __linux__ bool dma_buf_support = false; bool external_memory_fd_support = false; + bool external_semaphore_support = false; + bool external_semaphore_fd_support = false; #endif for (const auto& properties : ext_props) { @@ -6605,6 +6774,10 @@ static vk_device ggml_vk_get_device(size_t idx) { dma_buf_support = true; } else if (strcmp("VK_KHR_external_memory_fd", properties.extensionName) == 0) { external_memory_fd_support = true; + } else if (strcmp("VK_KHR_external_semaphore", properties.extensionName) == 0) { + external_semaphore_support = true; + } else if (strcmp("VK_KHR_external_semaphore_fd", properties.extensionName) == 0) { + external_semaphore_fd_support = true; #endif #if defined(VK_EXT_shader_64bit_indexing) } else if (strcmp("VK_EXT_shader_64bit_indexing", properties.extensionName) == 0) { @@ -6787,6 +6960,7 @@ static vk_device ggml_vk_get_device(size_t idx) { #ifdef __linux__ device->external_memory_dma_buf = dma_buf_support && external_memory_fd_support; + device->external_semaphore_fd = external_semaphore_support && external_semaphore_fd_support; #endif device->max_workgroup_size_log2 = uint32_t(log2f(float(device->properties.limits.maxComputeWorkGroupInvocations))); @@ -6967,6 +7141,10 @@ static vk_device ggml_vk_get_device(size_t idx) { device_extensions.push_back("VK_EXT_external_memory_dma_buf"); device_extensions.push_back("VK_KHR_external_memory_fd"); } + if (device->external_semaphore_fd) { + device_extensions.push_back("VK_KHR_external_semaphore"); + device_extensions.push_back("VK_KHR_external_semaphore_fd"); + } #endif #if defined(VK_EXT_shader_64bit_indexing) @@ -9057,6 +9235,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9073,6 +9252,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9094,6 +9274,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (src exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9111,6 +9292,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (dst exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9132,6 +9314,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: shared_staging\n", src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(buf_a); @@ -9178,6 +9361,14 @@ static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_devi return false; } + // Wait for any in-flight hop1 before destroying old buffers + if (path.hop1_fence_pending && path.hop1_device) { + VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), + "d2d grow wait hop1 fence"); + path.hop1_device->device.resetFences({ path.hop1_fence }); + path.hop1_fence_pending = false; + } + // Destroy old buffers ggml_vk_destroy_buffer(path.buf_a); ggml_vk_destroy_buffer(path.buf_b); @@ -9221,6 +9412,65 @@ static vk_d2d_path& ggml_vk_get_d2d_path(vk_device& src_dev, vk_device& dst_dev, return path; } + +static bool ggml_vk_d2d_is_async_capable(vk_device& src_dev, vk_device& dst_dev) { + std::lock_guard guard(vk_d2d_cache_mutex); + auto key = std::make_pair(src_dev.get(), dst_dev.get()); + auto it = vk_d2d_cache.find(key); + return it != vk_d2d_cache.end() && it->second.async_capable; +} + +static bool ggml_vk_buffer_copy_async_d2d( + vk_context& dst_compute_ctx, + vk_buffer& dst, size_t dst_offset, + vk_buffer& src, size_t src_offset, + size_t size) { + VK_LOG_DEBUG("ggml_vk_buffer_copy_async_d2d(" << size << ")"); + + vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); + + if (!path.async_capable || path.method == D2D_STAGING) { + return false; + } + + vk_buffer& src_side_buf = path.reverse_direction ? path.buf_b : path.buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? path.buf_a : path.buf_b; + + // Wait for any previous hop1 on this path to complete (command buffer reuse) + if (path.hop1_fence_pending) { + VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), + "d2d async wait hop1 fence"); + path.hop1_device->device.resetFences({ path.hop1_fence }); + path.hop1_fence_pending = false; + ggml_vk_command_pool_cleanup(src->device, path.hop1_cmd_pool); + } + + uint64_t signal_value = ++path.sem_value; + + // Hop 1: src device copies VRAM -> shared buffer, signals semaphore + { + std::lock_guard guard(src->device->mutex); + vk_context hop1_ctx = ggml_vk_create_temporary_context(path.hop1_cmd_pool); + ggml_vk_ctx_begin(src->device, hop1_ctx); + + VkBufferCopy bc{ src_offset, 0, size }; + vkCmdCopyBuffer(hop1_ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)src_side_buf->buffer, 1, &bc); + + hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, signal_value }); + + ggml_vk_ctx_end(hop1_ctx); + ggml_vk_submit(hop1_ctx, path.hop1_fence); + path.hop1_fence_pending = true; + } + + // Hop 2: record into dst compute context — waits on semaphore, copies shared buffer -> VRAM + dst_compute_ctx->s->wait_semaphores.push_back({ path.sem_dst, signal_value }); + + VkBufferCopy bc2{ 0, dst_offset, size }; + vkCmdCopyBuffer(dst_compute_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); + + return true; +} #endif static void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) { @@ -16985,8 +17235,19 @@ static bool ggml_backend_vk_cpy_tensor_async(ggml_backend_t backend_src, ggml_ba if (ggml_backend_buffer_is_vk(src->buffer)) { ggml_backend_vk_buffer_context * src_buf_ctx = (ggml_backend_vk_buffer_context *)src->buffer->context; - // Async copy only works within the same device if (src_buf_ctx->dev_buffer->device != dst_buf->device) { +#ifdef __linux__ + if (ggml_vk_d2d_is_async_capable(src_buf_ctx->dev_buffer->device, dst_buf->device)) { + vk_context compute_ctx = ggml_vk_get_compute_ctx(ctx); + if (ggml_vk_buffer_copy_async_d2d( + compute_ctx, + dst_buf, vk_tensor_offset(dst) + dst->view_offs, + src_buf_ctx->dev_buffer, vk_tensor_offset(src) + src->view_offs, + ggml_nbytes(src))) { + return true; + } + } +#endif return false; } From 3e53e38de359ee12d3708077193d2b240fb4f56d Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Tue, 26 May 2026 07:37:24 +0200 Subject: [PATCH 03/11] fixes, disable semaphore sharing on Nvidia + non-Nvidia --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 47 +++++++++++++++++++++------- 1 file changed, 36 insertions(+), 11 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index c8e4dff0b8..00294be3f2 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -2590,7 +2590,16 @@ vk_instance_t::~vk_instance_t() { { std::lock_guard guard(vk_d2d_cache_mutex); for (auto& entry : vk_d2d_cache) { - ggml_vk_d2d_destroy_shared_semaphore(entry.second); + // Neutralize entries without Vulkan API calls — device.destroy() + // in each device's destructor will implicitly free associated resources. + // Explicit Vulkan calls here are unsafe because the validation layer's + // static data may already be destroyed. + entry.second.async_capable = false; + entry.second.sem_src = VK_NULL_HANDLE; + entry.second.sem_dst = VK_NULL_HANDLE; + entry.second.hop1_fence = VK_NULL_HANDLE; + entry.second.hop1_fence_pending = false; + entry.second.hop1_cmd_pool.pool = nullptr; if (entry.second.host_ptr) { free(entry.second.host_ptr); entry.second.host_ptr = nullptr; @@ -4014,10 +4023,10 @@ static bool ggml_vk_d2d_check_timeline_semaphore_export(vk_device& dev) { ext_sem_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; ext_sem_info.pNext = &sem_type; - vk::ExternalSemaphoreProperties ext_sem_props = dev->physical_device.getExternalSemaphoreProperties(ext_sem_info); + vk::ExternalSemaphoreProperties props = dev->physical_device.getExternalSemaphoreProperties(ext_sem_info); - return (ext_sem_props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eExportable) && - (ext_sem_props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eImportable); + return (props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eExportable) && + (props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eImportable); } static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { @@ -4025,6 +4034,12 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d return false; } + bool src_nvidia = src_dev->vendor_id == VK_VENDOR_ID_NVIDIA; + bool dst_nvidia = dst_dev->vendor_id == VK_VENDOR_ID_NVIDIA; + if (src_nvidia != dst_nvidia) { + return false; + } + if (!ggml_vk_d2d_check_timeline_semaphore_export(src_dev) || !ggml_vk_d2d_check_timeline_semaphore_export(dst_dev)) { VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: timeline semaphore export/import not supported"); @@ -4034,9 +4049,7 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d vk::ExportSemaphoreCreateInfo export_ci; export_ci.handleTypes = vk::ExternalSemaphoreHandleTypeFlagBits::eOpaqueFd; - vk::SemaphoreTypeCreateInfo type_ci; - type_ci.semaphoreType = vk::SemaphoreType::eTimeline; - type_ci.initialValue = 0; + vk::SemaphoreTypeCreateInfo type_ci{ vk::SemaphoreType::eTimeline, 0 }; type_ci.pNext = &export_ci; vk::SemaphoreCreateInfo sem_ci; @@ -4099,14 +4112,24 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d } // fd ownership transferred to driver on successful import + try { + path.hop1_cmd_pool.init(src_dev, &src_dev->transfer_queue); + path.hop1_fence = src_dev->device.createFence({}); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: cmd pool/fence creation failed: " << e.what()); + if (path.hop1_cmd_pool.pool) { + path.hop1_cmd_pool.destroy(src_dev->device); + } + dst_dev->device.destroySemaphore(dst_sem); + src_dev->device.destroySemaphore(src_sem); + return false; + } + path.sem_src = src_sem; path.sem_dst = dst_sem; path.sem_value = 0; path.sem_src_device = src_dev.get(); path.sem_dst_device = dst_dev.get(); - - path.hop1_cmd_pool.init(src_dev, &src_dev->transfer_queue); - path.hop1_fence = src_dev->device.createFence({}); path.hop1_fence_pending = false; path.hop1_device = src_dev.get(); @@ -9463,7 +9486,9 @@ static bool ggml_vk_buffer_copy_async_d2d( path.hop1_fence_pending = true; } - // Hop 2: record into dst compute context — waits on semaphore, copies shared buffer -> VRAM + // Hop 2: start a new submission in the dst compute context so this copy + // waits only for its own hop1, not for later hop1s that overwrite the shared buffer. + ggml_vk_ctx_begin(dst->device, dst_compute_ctx); dst_compute_ctx->s->wait_semaphores.push_back({ path.sem_dst, signal_value }); VkBufferCopy bc2{ 0, dst_offset, size }; From 4f4cea311265ff49767c3525df473c50507f266e Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Tue, 26 May 2026 08:05:14 +0200 Subject: [PATCH 04/11] double buffering --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 181 +++++++++++++++++---------- 1 file changed, 118 insertions(+), 63 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 00294be3f2..eef33d93fd 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1215,11 +1215,21 @@ enum vk_d2d_method { struct vk_d2d_path { vk_d2d_method method = D2D_UNTESTED; bool reverse_direction = false; - vk_buffer buf_a; - vk_buffer buf_b; - void * host_ptr = nullptr; size_t size = 0; + static constexpr size_t POOL_SIZE = 2; + + struct slot { + vk_buffer buf_a; + vk_buffer buf_b; + void * host_ptr = nullptr; + uint64_t hop2_done = 0; + }; + + slot slots[POOL_SIZE]; + size_t num_slots = 0; + size_t pool_idx = 0; + bool async_capable = false; vk::Semaphore sem_src = VK_NULL_HANDLE; vk::Semaphore sem_dst = VK_NULL_HANDLE; @@ -2583,6 +2593,8 @@ static vk_instance_t vk_instance; #ifdef __linux__ static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path); +static bool ggml_vk_d2d_grow_slot(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, + size_t needed, vk_d2d_path::slot& s); #endif vk_instance_t::~vk_instance_t() { @@ -2600,15 +2612,17 @@ vk_instance_t::~vk_instance_t() { entry.second.hop1_fence = VK_NULL_HANDLE; entry.second.hop1_fence_pending = false; entry.second.hop1_cmd_pool.pool = nullptr; - if (entry.second.host_ptr) { - free(entry.second.host_ptr); - entry.second.host_ptr = nullptr; - } - if (entry.second.buf_a) { - entry.second.buf_a->size = 0; - } - if (entry.second.buf_b) { - entry.second.buf_b->size = 0; + for (auto& s : entry.second.slots) { + if (s.host_ptr) { + free(s.host_ptr); + s.host_ptr = nullptr; + } + if (s.buf_a) { + s.buf_a->size = 0; + } + if (s.buf_b) { + s.buf_b->size = 0; + } } } vk_d2d_cache.clear(); @@ -2625,15 +2639,17 @@ vk_device_struct::~vk_device_struct() { for (auto it = vk_d2d_cache.begin(); it != vk_d2d_cache.end(); ) { if (it->first.first == this || it->first.second == this) { ggml_vk_d2d_destroy_shared_semaphore(it->second); - if (it->second.host_ptr) { - free(it->second.host_ptr); - it->second.host_ptr = nullptr; - } - if (it->second.buf_a) { - it->second.buf_a->size = 0; - } - if (it->second.buf_b) { - it->second.buf_b->size = 0; + for (auto& s : it->second.slots) { + if (s.host_ptr) { + free(s.host_ptr); + s.host_ptr = nullptr; + } + if (s.buf_a) { + s.buf_a->size = 0; + } + if (s.buf_b) { + s.buf_b->size = 0; + } } it = vk_d2d_cache.erase(it); } else { @@ -4135,8 +4151,16 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d path.async_capable = true; - GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: async semaphore created\n", - src_dev->name.c_str(), dst_dev->name.c_str()); + // Allocate additional pool slots for double buffering + for (size_t i = path.num_slots; i < vk_d2d_path::POOL_SIZE; i++) { + if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, path.size, path.slots[i])) { + break; + } + path.num_slots = i + 1; + } + + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: async semaphore created (%zu pool slots)\n", + src_dev->name.c_str(), dst_dev->name.c_str(), path.num_slots); return true; } #endif @@ -9253,8 +9277,9 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev if (ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { path.method = D2D_DMABUF_P2P; path.reverse_direction = false; - path.buf_a = exp_buf; - path.buf_b = imp_buf; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); @@ -9270,8 +9295,9 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev if (ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { path.method = D2D_DMABUF_P2P; path.reverse_direction = true; - path.buf_a = exp_buf; - path.buf_b = imp_buf; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); @@ -9292,8 +9318,9 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { path.method = D2D_DMABUF_GTT; path.reverse_direction = false; - path.buf_a = exp_buf; - path.buf_b = imp_buf; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (src exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); @@ -9310,8 +9337,9 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { path.method = D2D_DMABUF_GTT; path.reverse_direction = true; - path.buf_a = exp_buf; - path.buf_b = imp_buf; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (dst exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); @@ -9331,9 +9359,10 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev if (ggml_vk_d2d_test_copy(src_dev, buf_a, VK_D2D_PROBE_SIZE) && ggml_vk_d2d_test_copy(dst_dev, buf_b, VK_D2D_PROBE_SIZE)) { path.method = D2D_SHARED_STAGING; - path.buf_a = buf_a; - path.buf_b = buf_b; - path.host_ptr = host_ptr; + path.slots[0].buf_a = buf_a; + path.slots[0].buf_b = buf_b; + path.slots[0].host_ptr = host_ptr; + path.num_slots = 1; path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: shared_staging\n", src_dev->name.c_str(), dst_dev->name.c_str()); @@ -9353,9 +9382,8 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev return path; } -static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, size_t needed) { - VK_LOG_DEBUG("ggml_vk_d2d_grow_path(" << needed << ", current=" << path.size << ")"); - +static bool ggml_vk_d2d_grow_slot(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, + size_t needed, vk_d2d_path::slot& s) { vk_buffer new_buf_a, new_buf_b; void * new_host_ptr = nullptr; bool ok = false; @@ -9384,6 +9412,21 @@ static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_devi return false; } + ggml_vk_destroy_buffer(s.buf_a); + ggml_vk_destroy_buffer(s.buf_b); + if (s.host_ptr) { + free(s.host_ptr); + } + + s.buf_a = new_buf_a; + s.buf_b = new_buf_b; + s.host_ptr = new_host_ptr; + return true; +} + +static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, size_t needed) { + VK_LOG_DEBUG("ggml_vk_d2d_grow_path(" << needed << ", current=" << path.size << ")"); + // Wait for any in-flight hop1 before destroying old buffers if (path.hop1_fence_pending && path.hop1_device) { VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), @@ -9392,16 +9435,12 @@ static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_devi path.hop1_fence_pending = false; } - // Destroy old buffers - ggml_vk_destroy_buffer(path.buf_a); - ggml_vk_destroy_buffer(path.buf_b); - if (path.host_ptr) { - free(path.host_ptr); + for (size_t i = 0; i < path.num_slots; i++) { + if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, needed, path.slots[i])) { + return false; + } } - path.buf_a = new_buf_a; - path.buf_b = new_buf_b; - path.host_ptr = new_host_ptr; path.size = needed; return true; } @@ -9422,12 +9461,15 @@ static vk_d2d_path& ggml_vk_get_d2d_path(vk_device& src_dev, vk_device& dst_dev, if (!ggml_vk_d2d_grow_path(path, src_dev, dst_dev, size)) { GGML_LOG_WARN("ggml_vulkan: d2d grow failed for %s -> %s, falling back to staging\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_destroy_buffer(path.buf_a); - ggml_vk_destroy_buffer(path.buf_b); - if (path.host_ptr) { - free(path.host_ptr); - path.host_ptr = nullptr; + for (size_t i = 0; i < path.num_slots; i++) { + ggml_vk_destroy_buffer(path.slots[i].buf_a); + ggml_vk_destroy_buffer(path.slots[i].buf_b); + if (path.slots[i].host_ptr) { + free(path.slots[i].host_ptr); + path.slots[i].host_ptr = nullptr; + } } + path.num_slots = 0; path.method = D2D_STAGING; path.size = 0; } @@ -9456,9 +9498,6 @@ static bool ggml_vk_buffer_copy_async_d2d( return false; } - vk_buffer& src_side_buf = path.reverse_direction ? path.buf_b : path.buf_a; - vk_buffer& dst_side_buf = path.reverse_direction ? path.buf_a : path.buf_b; - // Wait for any previous hop1 on this path to complete (command buffer reuse) if (path.hop1_fence_pending) { VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), @@ -9468,32 +9507,50 @@ static bool ggml_vk_buffer_copy_async_d2d( ggml_vk_command_pool_cleanup(src->device, path.hop1_cmd_pool); } - uint64_t signal_value = ++path.sem_value; + // Pick next pool slot (round-robin) + size_t slot_idx = path.pool_idx; + path.pool_idx = (path.pool_idx + 1) % path.num_slots; + vk_d2d_path::slot& slot = path.slots[slot_idx]; - // Hop 1: src device copies VRAM -> shared buffer, signals semaphore + vk_buffer& src_side_buf = path.reverse_direction ? slot.buf_b : slot.buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? slot.buf_a : slot.buf_b; + + // Two sem values per copy: hop1_signal for hop1→hop2, hop2_signal for hop2→next reuse + uint64_t hop1_signal = path.sem_value + 1; + uint64_t hop2_signal = path.sem_value + 2; + path.sem_value = hop2_signal; + + // Hop 1: src device copies VRAM -> shared buffer + // Wait for previous hop2 on this slot to finish reading before overwriting { std::lock_guard guard(src->device->mutex); vk_context hop1_ctx = ggml_vk_create_temporary_context(path.hop1_cmd_pool); ggml_vk_ctx_begin(src->device, hop1_ctx); + if (slot.hop2_done > 0) { + hop1_ctx->s->wait_semaphores.push_back({ path.sem_src, slot.hop2_done }); + } + VkBufferCopy bc{ src_offset, 0, size }; vkCmdCopyBuffer(hop1_ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)src_side_buf->buffer, 1, &bc); - hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, signal_value }); + hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, hop1_signal }); ggml_vk_ctx_end(hop1_ctx); ggml_vk_submit(hop1_ctx, path.hop1_fence); path.hop1_fence_pending = true; } - // Hop 2: start a new submission in the dst compute context so this copy - // waits only for its own hop1, not for later hop1s that overwrite the shared buffer. + // Hop 2: new submission in dst compute context — wait for hop1, signal when done reading ggml_vk_ctx_begin(dst->device, dst_compute_ctx); - dst_compute_ctx->s->wait_semaphores.push_back({ path.sem_dst, signal_value }); + dst_compute_ctx->s->wait_semaphores.push_back({ path.sem_dst, hop1_signal }); VkBufferCopy bc2{ 0, dst_offset, size }; vkCmdCopyBuffer(dst_compute_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); + dst_compute_ctx->s->signal_semaphores.push_back({ path.sem_dst, hop2_signal }); + slot.hop2_done = hop2_signal; + return true; } #endif @@ -9529,12 +9586,10 @@ static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& sr if (path.method != D2D_STAGING) { // buf_a is on the src-side device, buf_b is on the dst-side device // For reverse_direction (dst exports), buf_a is on dst_dev and buf_b is on src_dev - vk_buffer& src_side_buf = path.reverse_direction ? path.buf_b : path.buf_a; - vk_buffer& dst_side_buf = path.reverse_direction ? path.buf_a : path.buf_b; + vk_buffer& src_side_buf = path.reverse_direction ? path.slots[0].buf_b : path.slots[0].buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? path.slots[0].buf_a : path.slots[0].buf_b; - // Hop 1: src GPU copies VRAM -> shared buffer (same-device copy on src) ggml_vk_buffer_copy(src_side_buf, 0, src, src_offset, size); - // Hop 2: dst GPU copies shared buffer -> VRAM (same-device copy on dst) ggml_vk_buffer_copy(dst, dst_offset, dst_side_buf, 0, size); return; } From 5f512ae71bb593632c2186100a5185f9af3fe812 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Tue, 26 May 2026 09:33:34 +0200 Subject: [PATCH 05/11] use sync_fd binary semaphores for cross-driver synchronization --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 455 ++++++++++++++++++++++----- 1 file changed, 379 insertions(+), 76 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index eef33d93fd..b804e10690 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1198,8 +1198,8 @@ struct vk_buffer_struct { } VK_LOG_DEBUG("~vk_buffer_struct(" << buffer << ", " << size << ")"); - device->device.freeMemory(device_memory); device->device.destroyBuffer(buffer); + device->device.freeMemory(device_memory); } }; @@ -1212,6 +1212,12 @@ enum vk_d2d_method { D2D_STAGING, }; +enum vk_d2d_sync_method { + D2D_SYNC_NONE, + D2D_SYNC_TIMELINE, + D2D_SYNC_SYNCFD, +}; + struct vk_d2d_path { vk_d2d_method method = D2D_UNTESTED; bool reverse_direction = false; @@ -1224,13 +1230,17 @@ struct vk_d2d_path { vk_buffer buf_b; void * host_ptr = nullptr; uint64_t hop2_done = 0; + // sync_fd back-edge: hop2 signals a per-copy semaphore (stored here), + // exported via sync_fd into a per-copy semaphore on src for next hop1 + vk::Semaphore last_back_sem = VK_NULL_HANDLE; + bool back_edge_ready = false; }; slot slots[POOL_SIZE]; size_t num_slots = 0; size_t pool_idx = 0; - bool async_capable = false; + vk_d2d_sync_method sync_method = D2D_SYNC_NONE; vk::Semaphore sem_src = VK_NULL_HANDLE; vk::Semaphore sem_dst = VK_NULL_HANDLE; uint64_t sem_value = 0; @@ -1238,9 +1248,10 @@ struct vk_d2d_path { vk_device_struct * sem_dst_device = nullptr; vk_command_pool hop1_cmd_pool; - vk::Fence hop1_fence = VK_NULL_HANDLE; - bool hop1_fence_pending = false; vk_device_struct * hop1_device = nullptr; + + vk_command_pool hop2_cmd_pool; + vk_device_struct * hop2_device = nullptr; }; #endif @@ -2592,7 +2603,7 @@ static bool vk_instance_initialized = false; static vk_instance_t vk_instance; #ifdef __linux__ -static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path); +static void ggml_vk_d2d_destroy_sync(vk_d2d_path& path); static bool ggml_vk_d2d_grow_slot(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, size_t needed, vk_d2d_path::slot& s); #endif @@ -2606,12 +2617,11 @@ vk_instance_t::~vk_instance_t() { // in each device's destructor will implicitly free associated resources. // Explicit Vulkan calls here are unsafe because the validation layer's // static data may already be destroyed. - entry.second.async_capable = false; + entry.second.sync_method = D2D_SYNC_NONE; entry.second.sem_src = VK_NULL_HANDLE; entry.second.sem_dst = VK_NULL_HANDLE; - entry.second.hop1_fence = VK_NULL_HANDLE; - entry.second.hop1_fence_pending = false; entry.second.hop1_cmd_pool.pool = nullptr; + entry.second.hop2_cmd_pool.pool = nullptr; for (auto& s : entry.second.slots) { if (s.host_ptr) { free(s.host_ptr); @@ -2623,6 +2633,8 @@ vk_instance_t::~vk_instance_t() { if (s.buf_b) { s.buf_b->size = 0; } + s.last_back_sem = VK_NULL_HANDLE; + s.back_edge_ready = false; } } vk_d2d_cache.clear(); @@ -2638,7 +2650,7 @@ vk_device_struct::~vk_device_struct() { std::lock_guard guard(vk_d2d_cache_mutex); for (auto it = vk_d2d_cache.begin(); it != vk_d2d_cache.end(); ) { if (it->first.first == this || it->first.second == this) { - ggml_vk_d2d_destroy_shared_semaphore(it->second); + ggml_vk_d2d_destroy_sync(it->second); for (auto& s : it->second.slots) { if (s.host_ptr) { free(s.host_ptr); @@ -3456,9 +3468,17 @@ static vk_context ggml_vk_create_temporary_context(vk_command_pool& p) { return result; } -static vk_semaphore * ggml_vk_create_binary_semaphore(ggml_backend_vk_context * ctx) { - VK_LOG_DEBUG("ggml_vk_create_timeline_semaphore()"); +static vk_semaphore * ggml_vk_create_binary_semaphore(ggml_backend_vk_context * ctx, + vk::ExternalSemaphoreHandleTypeFlags export_handle_types = {}) { + VK_LOG_DEBUG("ggml_vk_create_binary_semaphore()"); + vk::ExportSemaphoreCreateInfo export_ci; + export_ci.handleTypes = export_handle_types; + vk::SemaphoreTypeCreateInfo tci{ vk::SemaphoreType::eBinary, 0 }; + if (export_handle_types) { + tci.pNext = &export_ci; + } + vk::SemaphoreCreateInfo ci{}; ci.setPNext(&tci); vk::Semaphore semaphore = ctx->device->device.createSemaphore(ci); @@ -3994,26 +4014,39 @@ static bool ggml_vk_d2d_try_shared_staging(vk_device& dev_a, vk_device& dev_b, s return true; } -static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path) { - if (!path.async_capable) { +static void ggml_vk_d2d_destroy_sync(vk_d2d_path& path) { + if (path.sync_method == D2D_SYNC_NONE) { return; } - if (path.hop1_fence_pending && path.hop1_device) { - try { - VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), - "d2d destroy wait hop1 fence"); - } catch (...) {} - path.hop1_fence_pending = false; + if (path.sync_method == D2D_SYNC_TIMELINE) { + if (path.sem_value > 0 && path.sem_src_device) { + try { + VkSemaphoreWaitInfo wait_info = {}; + wait_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO; + VkSemaphore sem = path.sem_src; + uint64_t val = path.sem_value; + wait_info.semaphoreCount = 1; + wait_info.pSemaphores = &sem; + wait_info.pValues = &val; + vkWaitSemaphores(path.sem_src_device->device, &wait_info, UINT64_MAX); + } catch (...) {} + } + } else if (path.sync_method == D2D_SYNC_SYNCFD) { + if (path.hop1_device) { + vkDeviceWaitIdle(path.hop1_device->device); + } + if (path.hop2_device) { + vkDeviceWaitIdle(path.hop2_device->device); + } } - if (path.hop1_fence) { - path.hop1_device->device.destroyFence(path.hop1_fence); - path.hop1_fence = VK_NULL_HANDLE; - } if (path.hop1_cmd_pool.pool) { path.hop1_cmd_pool.destroy(path.hop1_device->device); } + if (path.hop2_cmd_pool.pool) { + path.hop2_cmd_pool.destroy(path.hop2_device->device); + } if (path.sem_src) { path.sem_src_device->device.destroySemaphore(path.sem_src); @@ -4023,12 +4056,17 @@ static void ggml_vk_d2d_destroy_shared_semaphore(vk_d2d_path& path) { path.sem_dst_device->device.destroySemaphore(path.sem_dst); path.sem_dst = VK_NULL_HANDLE; } + for (size_t i = 0; i < path.num_slots; i++) { + path.slots[i].last_back_sem = VK_NULL_HANDLE; + path.slots[i].back_edge_ready = false; + } - path.async_capable = false; + path.sync_method = D2D_SYNC_NONE; path.sem_value = 0; path.sem_src_device = nullptr; path.sem_dst_device = nullptr; path.hop1_device = nullptr; + path.hop2_device = nullptr; } static bool ggml_vk_d2d_check_timeline_semaphore_export(vk_device& dev) { @@ -4045,20 +4083,32 @@ static bool ggml_vk_d2d_check_timeline_semaphore_export(vk_device& dev) { (props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eImportable); } -static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { +static bool ggml_vk_d2d_check_sync_fd_support(vk_device& dev) { + if (!dev->external_semaphore_fd) { + return false; + } + + vk::PhysicalDeviceExternalSemaphoreInfo ext_sem_info; + ext_sem_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd; + + vk::ExternalSemaphoreProperties props = dev->physical_device.getExternalSemaphoreProperties(ext_sem_info); + + return (props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eExportable) && + (props.externalSemaphoreFeatures & vk::ExternalSemaphoreFeatureFlagBits::eImportable); +} + +static bool ggml_vk_d2d_try_timeline_sync(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { if (!src_dev->external_semaphore_fd || !dst_dev->external_semaphore_fd) { return false; } - bool src_nvidia = src_dev->vendor_id == VK_VENDOR_ID_NVIDIA; - bool dst_nvidia = dst_dev->vendor_id == VK_VENDOR_ID_NVIDIA; - if (src_nvidia != dst_nvidia) { + if (src_dev->driver_id != dst_dev->driver_id) { return false; } if (!ggml_vk_d2d_check_timeline_semaphore_export(src_dev) || !ggml_vk_d2d_check_timeline_semaphore_export(dst_dev)) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: timeline semaphore export/import not supported"); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: timeline semaphore export/import not supported"); return false; } @@ -4075,7 +4125,7 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d try { src_sem = src_dev->device.createSemaphore(sem_ci); } catch (const vk::SystemError& e) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: createSemaphore on src failed: " << e.what()); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: createSemaphore on src failed: " << e.what()); return false; } @@ -4087,12 +4137,12 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d try { fd = src_dev->device.getSemaphoreFdKHR(get_fd_info); } catch (const vk::SystemError& e) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: getSemaphoreFdKHR failed: " << e.what()); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: getSemaphoreFdKHR failed: " << e.what()); src_dev->device.destroySemaphore(src_sem); return false; } if (fd < 0) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: getSemaphoreFdKHR returned invalid fd"); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: getSemaphoreFdKHR returned invalid fd"); src_dev->device.destroySemaphore(src_sem); return false; } @@ -4106,7 +4156,7 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d try { dst_sem = dst_dev->device.createSemaphore(dst_sem_ci); } catch (const vk::SystemError& e) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: createSemaphore on dst failed: " << e.what()); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: createSemaphore on dst failed: " << e.what()); close(fd); src_dev->device.destroySemaphore(src_sem); return false; @@ -4120,19 +4170,17 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d try { dst_dev->device.importSemaphoreFdKHR(import_info); } catch (const vk::SystemError& e) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: importSemaphoreFdKHR failed: " << e.what()); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: importSemaphoreFdKHR failed: " << e.what()); close(fd); dst_dev->device.destroySemaphore(dst_sem); src_dev->device.destroySemaphore(src_sem); return false; } - // fd ownership transferred to driver on successful import try { path.hop1_cmd_pool.init(src_dev, &src_dev->transfer_queue); - path.hop1_fence = src_dev->device.createFence({}); } catch (const vk::SystemError& e) { - VK_LOG_DEBUG("ggml_vk_d2d_create_shared_semaphore: cmd pool/fence creation failed: " << e.what()); + VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: cmd pool/fence creation failed: " << e.what()); if (path.hop1_cmd_pool.pool) { path.hop1_cmd_pool.destroy(src_dev->device); } @@ -4146,12 +4194,9 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d path.sem_value = 0; path.sem_src_device = src_dev.get(); path.sem_dst_device = dst_dev.get(); - path.hop1_fence_pending = false; path.hop1_device = src_dev.get(); + path.sync_method = D2D_SYNC_TIMELINE; - path.async_capable = true; - - // Allocate additional pool slots for double buffering for (size_t i = path.num_slots; i < vk_d2d_path::POOL_SIZE; i++) { if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, path.size, path.slots[i])) { break; @@ -4159,10 +4204,80 @@ static bool ggml_vk_d2d_create_shared_semaphore(vk_device& src_dev, vk_device& d path.num_slots = i + 1; } - GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: async semaphore created (%zu pool slots)\n", + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: timeline sync (same driver), %zu pool slots\n", src_dev->name.c_str(), dst_dev->name.c_str(), path.num_slots); return true; } + +static bool ggml_vk_d2d_try_syncfd_sync(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { + if (!ggml_vk_d2d_check_sync_fd_support(src_dev) || + !ggml_vk_d2d_check_sync_fd_support(dst_dev)) { + VK_LOG_DEBUG("ggml_vk_d2d_try_syncfd_sync: sync_fd not supported on one or both devices"); + return false; + } + + vk::ExportSemaphoreCreateInfo export_ci; + export_ci.handleTypes = vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd; + + vk::SemaphoreCreateInfo src_sem_ci; + src_sem_ci.pNext = &export_ci; + + vk::Semaphore src_sem; + try { + src_sem = src_dev->device.createSemaphore(src_sem_ci); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_try_syncfd_sync: createSemaphore on src failed: " << e.what()); + return false; + } + + try { + path.hop1_cmd_pool.init(src_dev, &src_dev->compute_queue); + path.hop2_cmd_pool.init(dst_dev, &dst_dev->compute_queue); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_try_syncfd_sync: cmd pool creation failed: " << e.what()); + if (path.hop2_cmd_pool.pool) { + path.hop2_cmd_pool.destroy(dst_dev->device); + } + if (path.hop1_cmd_pool.pool) { + path.hop1_cmd_pool.destroy(src_dev->device); + } + src_dev->device.destroySemaphore(src_sem); + return false; + } + + path.sem_src = src_sem; + path.sem_value = 0; + path.sem_src_device = src_dev.get(); + path.sem_dst_device = dst_dev.get(); + path.hop1_device = src_dev.get(); + path.hop2_device = dst_dev.get(); + path.sync_method = D2D_SYNC_SYNCFD; + + // Grow to double-buffer pool size + for (size_t i = path.num_slots; i < vk_d2d_path::POOL_SIZE; i++) { + if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, path.size, path.slots[i])) { + break; + } + path.num_slots = i + 1; + } + + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: sync_fd sync (cross-driver), %zu pool slots\n", + src_dev->name.c_str(), dst_dev->name.c_str(), path.num_slots); + return true; +} + +static void ggml_vk_d2d_setup_sync(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path) { + if (ggml_vk_d2d_try_timeline_sync(src_dev, dst_dev, path)) { + return; + } + + if (ggml_vk_d2d_try_syncfd_sync(src_dev, dst_dev, path)) { + return; + } + + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: no async sync available, using CPU fallback\n", + src_dev->name.c_str(), dst_dev->name.c_str()); +} #endif static vk_subbuffer ggml_vk_subbuffer(const ggml_backend_vk_context* ctx, const vk_buffer& buf, size_t offset = 0) { @@ -9283,7 +9398,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9301,7 +9416,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9324,7 +9439,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (src exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9343,7 +9458,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_gtt (dst exports GTT)\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(exp_buf); @@ -9366,7 +9481,7 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev path.size = VK_D2D_PROBE_SIZE; GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: shared_staging\n", src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_create_shared_semaphore(src_dev, dst_dev, path); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); return path; } ggml_vk_destroy_buffer(buf_a); @@ -9427,18 +9542,43 @@ static bool ggml_vk_d2d_grow_slot(vk_d2d_path& path, vk_device& src_dev, vk_devi static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_device& dst_dev, size_t needed) { VK_LOG_DEBUG("ggml_vk_d2d_grow_path(" << needed << ", current=" << path.size << ")"); - // Wait for any in-flight hop1 before destroying old buffers - if (path.hop1_fence_pending && path.hop1_device) { - VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), - "d2d grow wait hop1 fence"); - path.hop1_device->device.resetFences({ path.hop1_fence }); - path.hop1_fence_pending = false; + // Wait for all in-flight work on both devices, then release all command buffer + // references to the shared buffers before destroying them. + // The shared buffers may be referenced by: hop1_cmd_pool command buffers, + // device queue command buffers (from probe test copies and sync d2d copies), + // and compute context command buffers (from hop2). + if (path.sync_method == D2D_SYNC_TIMELINE && path.sem_value > 0) { + VkSemaphoreWaitInfo wait_info = {}; + wait_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO; + VkSemaphore sem = path.sem_src; + uint64_t val = path.sem_value; + wait_info.semaphoreCount = 1; + wait_info.pSemaphores = &sem; + wait_info.pValues = &val; + vkWaitSemaphores(path.sem_src_device->device, &wait_info, UINT64_MAX); + + path.sem_src_device->device.resetCommandPool(path.hop1_cmd_pool.pool); + for (auto& cb : path.hop1_cmd_pool.cmd_buffers) { + cb.in_use = false; + } + } else if (path.sync_method == D2D_SYNC_SYNCFD) { + // No fences to wait — back-edge semaphores are GPU-only. + // deviceWaitIdle below ensures all work is done. } + src_dev->device.waitIdle(); + dst_dev->device.waitIdle(); + for (size_t i = 0; i < path.num_slots; i++) { if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, needed, path.slots[i])) { return false; } + path.slots[i].back_edge_ready = false; + } + + if (path.sync_method == D2D_SYNC_SYNCFD) { + ggml_vk_command_pool_cleanup(src_dev, path.hop1_cmd_pool); + ggml_vk_command_pool_cleanup(dst_dev, path.hop2_cmd_pool); } path.size = needed; @@ -9482,28 +9622,26 @@ static bool ggml_vk_d2d_is_async_capable(vk_device& src_dev, vk_device& dst_dev) std::lock_guard guard(vk_d2d_cache_mutex); auto key = std::make_pair(src_dev.get(), dst_dev.get()); auto it = vk_d2d_cache.find(key); - return it != vk_d2d_cache.end() && it->second.async_capable; + return it != vk_d2d_cache.end() && it->second.sync_method != D2D_SYNC_NONE; } -static bool ggml_vk_buffer_copy_async_d2d( +static bool ggml_vk_buffer_copy_async_d2d_timeline( vk_context& dst_compute_ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, - size_t size) { - VK_LOG_DEBUG("ggml_vk_buffer_copy_async_d2d(" << size << ")"); + size_t size, + vk_d2d_path& path) { - vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); - - if (!path.async_capable || path.method == D2D_STAGING) { - return false; - } - - // Wait for any previous hop1 on this path to complete (command buffer reuse) - if (path.hop1_fence_pending) { - VK_CHECK(path.hop1_device->device.waitForFences({ path.hop1_fence }, true, UINT64_MAX), - "d2d async wait hop1 fence"); - path.hop1_device->device.resetFences({ path.hop1_fence }); - path.hop1_fence_pending = false; + // Periodic command pool cleanup + if (path.hop1_cmd_pool.buffers_in_use() >= 8) { + VkSemaphoreWaitInfo wait_info = {}; + wait_info.sType = VK_STRUCTURE_TYPE_SEMAPHORE_WAIT_INFO; + VkSemaphore sem = path.sem_src; + uint64_t val = path.sem_value; + wait_info.semaphoreCount = 1; + wait_info.pSemaphores = &sem; + wait_info.pValues = &val; + vkWaitSemaphores(path.sem_src_device->device, &wait_info, UINT64_MAX); ggml_vk_command_pool_cleanup(src->device, path.hop1_cmd_pool); } @@ -9515,13 +9653,11 @@ static bool ggml_vk_buffer_copy_async_d2d( vk_buffer& src_side_buf = path.reverse_direction ? slot.buf_b : slot.buf_a; vk_buffer& dst_side_buf = path.reverse_direction ? slot.buf_a : slot.buf_b; - // Two sem values per copy: hop1_signal for hop1→hop2, hop2_signal for hop2→next reuse uint64_t hop1_signal = path.sem_value + 1; uint64_t hop2_signal = path.sem_value + 2; path.sem_value = hop2_signal; // Hop 1: src device copies VRAM -> shared buffer - // Wait for previous hop2 on this slot to finish reading before overwriting { std::lock_guard guard(src->device->mutex); vk_context hop1_ctx = ggml_vk_create_temporary_context(path.hop1_cmd_pool); @@ -9537,11 +9673,10 @@ static bool ggml_vk_buffer_copy_async_d2d( hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, hop1_signal }); ggml_vk_ctx_end(hop1_ctx); - ggml_vk_submit(hop1_ctx, path.hop1_fence); - path.hop1_fence_pending = true; + ggml_vk_submit(hop1_ctx, {}); } - // Hop 2: new submission in dst compute context — wait for hop1, signal when done reading + // Hop 2: deferred in dst compute context ggml_vk_ctx_begin(dst->device, dst_compute_ctx); dst_compute_ctx->s->wait_semaphores.push_back({ path.sem_dst, hop1_signal }); @@ -9553,6 +9688,140 @@ static bool ggml_vk_buffer_copy_async_d2d( return true; } + +static bool ggml_vk_d2d_syncfd_export_import(vk_device& from_dev, vk::Semaphore from_sem, + vk_device& to_dev, vk::Semaphore to_sem) { + vk::SemaphoreGetFdInfoKHR get_fd_info; + get_fd_info.semaphore = from_sem; + get_fd_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd; + + int fd; + try { + fd = from_dev->device.getSemaphoreFdKHR(get_fd_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_syncfd_export_import: getSemaphoreFdKHR failed: " << e.what()); + return false; + } + + vk::ImportSemaphoreFdInfoKHR import_info; + import_info.semaphore = to_sem; + import_info.handleType = vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd; + import_info.fd = fd; + import_info.flags = vk::SemaphoreImportFlagBits::eTemporary; + + try { + to_dev->device.importSemaphoreFdKHR(import_info); + } catch (const vk::SystemError& e) { + VK_LOG_DEBUG("ggml_vk_d2d_syncfd_export_import: importSemaphoreFdKHR failed: " << e.what()); + close(fd); + return false; + } + + return true; +} + +static bool ggml_vk_buffer_copy_async_d2d_syncfd( + ggml_backend_vk_context * src_ctx, + ggml_backend_vk_context * dst_ctx, + vk_buffer& dst, size_t dst_offset, + vk_buffer& src, size_t src_offset, + size_t size, + vk_d2d_path& path) { + + size_t slot_idx = path.pool_idx; + path.pool_idx = (path.pool_idx + 1) % path.num_slots; + vk_d2d_path::slot& slot = path.slots[slot_idx]; + + vk_buffer& src_side_buf = path.reverse_direction ? slot.buf_b : slot.buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? slot.buf_a : slot.buf_b; + + // Create all per-copy semaphores upfront to avoid vector reallocation + // invalidating pointers (ggml_vk_create_binary_semaphore returns pointer + // into a vector that may reallocate on subsequent push_back) + vk::Semaphore back_wait_sem = VK_NULL_HANDLE; + if (slot.back_edge_ready) { + back_wait_sem = ggml_vk_create_binary_semaphore(src_ctx)->s; + slot.back_edge_ready = false; + } + vk::Semaphore fwd_sem = ggml_vk_create_binary_semaphore(dst_ctx)->s; + vk::Semaphore back_signal_sem = ggml_vk_create_binary_semaphore(dst_ctx, + vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd)->s; + + // Back edge: export previous hop2's signal into the src-side wait semaphore + if (back_wait_sem) { + if (!ggml_vk_d2d_syncfd_export_import(dst->device, slot.last_back_sem, src->device, back_wait_sem)) { + return false; + } + } + + // Hop 1: src device copies VRAM -> shared buffer + { + std::lock_guard guard(src->device->mutex); + vk_context hop1_ctx = ggml_vk_create_temporary_context(path.hop1_cmd_pool); + ggml_vk_ctx_begin(src->device, hop1_ctx); + + if (back_wait_sem) { + hop1_ctx->s->wait_semaphores.push_back({ back_wait_sem, 0 }); + } + + VkBufferCopy bc{ src_offset, 0, size }; + vkCmdCopyBuffer(hop1_ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)src_side_buf->buffer, 1, &bc); + + hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, 0 }); + + ggml_vk_ctx_end(hop1_ctx); + ggml_vk_submit(hop1_ctx, {}); + } + + // Forward edge: export sync_fd from sem_src, import into per-copy dst semaphore + if (!ggml_vk_d2d_syncfd_export_import(src->device, path.sem_src, dst->device, fwd_sem)) { + return false; + } + + // Hop 2: dst device copies shared buffer -> VRAM + { + std::lock_guard guard(dst->device->mutex); + vk_context hop2_ctx = ggml_vk_create_temporary_context(path.hop2_cmd_pool); + ggml_vk_ctx_begin(dst->device, hop2_ctx); + + hop2_ctx->s->wait_semaphores.push_back({ fwd_sem, 0 }); + + VkBufferCopy bc2{ 0, dst_offset, size }; + vkCmdCopyBuffer(hop2_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); + + hop2_ctx->s->signal_semaphores.push_back({ back_signal_sem, 0 }); + + ggml_vk_ctx_end(hop2_ctx); + ggml_vk_submit(hop2_ctx, {}); + } + + slot.last_back_sem = back_signal_sem; + slot.back_edge_ready = true; + + return true; +} + +static bool ggml_vk_buffer_copy_async_d2d( + ggml_backend_vk_context * src_ctx, + ggml_backend_vk_context * dst_ctx, + vk_buffer& dst, size_t dst_offset, + vk_buffer& src, size_t src_offset, + size_t size) { + VK_LOG_DEBUG("ggml_vk_buffer_copy_async_d2d(" << size << ")"); + + vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); + + if (path.sync_method == D2D_SYNC_NONE || path.method == D2D_STAGING) { + return false; + } + + if (path.sync_method == D2D_SYNC_TIMELINE) { + vk_context compute_ctx = ggml_vk_get_compute_ctx(dst_ctx); + return ggml_vk_buffer_copy_async_d2d_timeline(compute_ctx, dst, dst_offset, src, src_offset, size, path); + } + + return ggml_vk_buffer_copy_async_d2d_syncfd(src_ctx, dst_ctx, dst, dst_offset, src, src_offset, size, path); +} #endif static void ggml_vk_buffer_copy_async(vk_context& ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size) { @@ -16810,6 +17079,24 @@ static void ggml_vk_graph_cleanup(ggml_backend_vk_context * ctx) { } ctx->gc.semaphores.clear(); +#ifdef __linux__ + { + std::lock_guard guard(vk_d2d_cache_mutex); + for (auto& entry : vk_d2d_cache) { + vk_d2d_path& path = entry.second; + if (path.sync_method != D2D_SYNC_SYNCFD) { + continue; + } + if (path.sem_dst_device == ctx->device.get() || path.sem_src_device == ctx->device.get()) { + for (size_t i = 0; i < path.num_slots; i++) { + path.slots[i].last_back_sem = VK_NULL_HANDLE; + path.slots[i].back_edge_ready = false; + } + } + } + } +#endif + for (size_t i = 0; i < ctx->gc.tl_semaphores.size(); i++) { ctx->device->device.destroySemaphore({ ctx->gc.tl_semaphores[i].s }); } @@ -16838,6 +17125,22 @@ static void ggml_vk_cleanup(ggml_backend_vk_context * ctx) { ggml_vk_graph_cleanup(ctx); +#ifdef __linux__ + { + vkDeviceWaitIdle(ctx->device->device); + std::lock_guard guard(vk_d2d_cache_mutex); + for (auto& entry : vk_d2d_cache) { + vk_d2d_path& path = entry.second; + if (path.hop1_device == ctx->device.get() && path.hop1_cmd_pool.pool) { + ggml_vk_command_pool_cleanup(ctx->device, path.hop1_cmd_pool); + } + if (path.hop2_device == ctx->device.get() && path.hop2_cmd_pool.pool) { + ggml_vk_command_pool_cleanup(ctx->device, path.hop2_cmd_pool); + } + } + } +#endif + ggml_vk_destroy_buffer(ctx->prealloc_x); ggml_vk_destroy_buffer(ctx->prealloc_y); ggml_vk_destroy_buffer(ctx->prealloc_split_k); @@ -17318,9 +17621,9 @@ static bool ggml_backend_vk_cpy_tensor_async(ggml_backend_t backend_src, ggml_ba if (src_buf_ctx->dev_buffer->device != dst_buf->device) { #ifdef __linux__ if (ggml_vk_d2d_is_async_capable(src_buf_ctx->dev_buffer->device, dst_buf->device)) { - vk_context compute_ctx = ggml_vk_get_compute_ctx(ctx); + ggml_backend_vk_context * src_ctx = (ggml_backend_vk_context *)backend_src->context; if (ggml_vk_buffer_copy_async_d2d( - compute_ctx, + src_ctx, ctx, dst_buf, vk_tensor_offset(dst) + dst->view_offs, src_buf_ctx->dev_buffer, vk_tensor_offset(src) + src->view_offs, ggml_nbytes(src))) { From 86ad44176517671051fb74f36778d66093c8b820 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Fri, 29 May 2026 15:08:39 +0200 Subject: [PATCH 06/11] use allocator instead of hop2 --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 140 ++++++++++++++------------- 1 file changed, 71 insertions(+), 69 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index b804e10690..e1bd9faac1 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1252,6 +1252,11 @@ struct vk_d2d_path { vk_command_pool hop2_cmd_pool; vk_device_struct * hop2_device = nullptr; + + // Syncfd bump sub-allocator: each copy claims alloc_offset space, + // reset between evaluations. high_water_mark drives buffer resizing. + size_t alloc_offset = 0; + size_t high_water_mark = 0; }; #endif @@ -4036,8 +4041,8 @@ static void ggml_vk_d2d_destroy_sync(vk_d2d_path& path) { if (path.hop1_device) { vkDeviceWaitIdle(path.hop1_device->device); } - if (path.hop2_device) { - vkDeviceWaitIdle(path.hop2_device->device); + if (path.sem_dst_device) { + vkDeviceWaitIdle(path.sem_dst_device->device); } } @@ -4056,13 +4061,11 @@ static void ggml_vk_d2d_destroy_sync(vk_d2d_path& path) { path.sem_dst_device->device.destroySemaphore(path.sem_dst); path.sem_dst = VK_NULL_HANDLE; } - for (size_t i = 0; i < path.num_slots; i++) { - path.slots[i].last_back_sem = VK_NULL_HANDLE; - path.slots[i].back_edge_ready = false; - } path.sync_method = D2D_SYNC_NONE; path.sem_value = 0; + path.alloc_offset = 0; + path.high_water_mark = 0; path.sem_src_device = nullptr; path.sem_dst_device = nullptr; path.hop1_device = nullptr; @@ -4232,12 +4235,8 @@ static bool ggml_vk_d2d_try_syncfd_sync(vk_device& src_dev, vk_device& dst_dev, try { path.hop1_cmd_pool.init(src_dev, &src_dev->compute_queue); - path.hop2_cmd_pool.init(dst_dev, &dst_dev->compute_queue); } catch (const vk::SystemError& e) { VK_LOG_DEBUG("ggml_vk_d2d_try_syncfd_sync: cmd pool creation failed: " << e.what()); - if (path.hop2_cmd_pool.pool) { - path.hop2_cmd_pool.destroy(dst_dev->device); - } if (path.hop1_cmd_pool.pool) { path.hop1_cmd_pool.destroy(src_dev->device); } @@ -4250,19 +4249,22 @@ static bool ggml_vk_d2d_try_syncfd_sync(vk_device& src_dev, vk_device& dst_dev, path.sem_src_device = src_dev.get(); path.sem_dst_device = dst_dev.get(); path.hop1_device = src_dev.get(); - path.hop2_device = dst_dev.get(); path.sync_method = D2D_SYNC_SYNCFD; + path.alloc_offset = 0; + path.high_water_mark = 0; - // Grow to double-buffer pool size - for (size_t i = path.num_slots; i < vk_d2d_path::POOL_SIZE; i++) { - if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, path.size, path.slots[i])) { - break; + // Single shared buffer pair for bump sub-allocator + if (path.num_slots == 0) { + if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, path.size, path.slots[0])) { + path.hop1_cmd_pool.destroy(src_dev->device); + src_dev->device.destroySemaphore(src_sem); + return false; } - path.num_slots = i + 1; + path.num_slots = 1; } - GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: sync_fd sync (cross-driver), %zu pool slots\n", - src_dev->name.c_str(), dst_dev->name.c_str(), path.num_slots); + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: sync_fd sync (cross-driver)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); return true; } @@ -9562,7 +9564,6 @@ static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_devi cb.in_use = false; } } else if (path.sync_method == D2D_SYNC_SYNCFD) { - // No fences to wait — back-edge semaphores are GPU-only. // deviceWaitIdle below ensures all work is done. } @@ -9573,12 +9574,14 @@ static bool ggml_vk_d2d_grow_path(vk_d2d_path& path, vk_device& src_dev, vk_devi if (!ggml_vk_d2d_grow_slot(path, src_dev, dst_dev, needed, path.slots[i])) { return false; } - path.slots[i].back_edge_ready = false; } if (path.sync_method == D2D_SYNC_SYNCFD) { ggml_vk_command_pool_cleanup(src_dev, path.hop1_cmd_pool); - ggml_vk_command_pool_cleanup(dst_dev, path.hop2_cmd_pool); + if (path.hop2_cmd_pool.pool) { + ggml_vk_command_pool_cleanup(dst_dev, path.hop2_cmd_pool); + } + path.alloc_offset = 0; } path.size = needed; @@ -9597,6 +9600,7 @@ static vk_d2d_path& ggml_vk_get_d2d_path(vk_device& src_dev, vk_device& dst_dev, vk_d2d_path& path = it->second; + // Grow buffer if a single copy doesn't fit if (path.method != D2D_STAGING && path.size < size) { if (!ggml_vk_d2d_grow_path(path, src_dev, dst_dev, size)) { GGML_LOG_WARN("ggml_vulkan: d2d grow failed for %s -> %s, falling back to staging\n", @@ -9615,6 +9619,14 @@ static vk_d2d_path& ggml_vk_get_d2d_path(vk_device& src_dev, vk_device& dst_dev, } } + // Syncfd: grow buffer to high-water mark at start of each evaluation + if (path.sync_method == D2D_SYNC_SYNCFD && + path.alloc_offset == 0 && path.high_water_mark > path.size) { + if (ggml_vk_d2d_grow_path(path, src_dev, dst_dev, path.high_water_mark)) { + path.high_water_mark = 0; + } + } + return path; } @@ -9721,50 +9733,48 @@ static bool ggml_vk_d2d_syncfd_export_import(vk_device& from_dev, vk::Semaphore } static bool ggml_vk_buffer_copy_async_d2d_syncfd( - ggml_backend_vk_context * src_ctx, ggml_backend_vk_context * dst_ctx, vk_buffer& dst, size_t dst_offset, vk_buffer& src, size_t src_offset, size_t size, vk_d2d_path& path) { - size_t slot_idx = path.pool_idx; - path.pool_idx = (path.pool_idx + 1) % path.num_slots; - vk_d2d_path::slot& slot = path.slots[slot_idx]; + static constexpr size_t D2D_BUMP_ALIGN = 256; + vk_d2d_path::slot& slot = path.slots[0]; vk_buffer& src_side_buf = path.reverse_direction ? slot.buf_b : slot.buf_a; vk_buffer& dst_side_buf = path.reverse_direction ? slot.buf_a : slot.buf_b; - // Create all per-copy semaphores upfront to avoid vector reallocation - // invalidating pointers (ggml_vk_create_binary_semaphore returns pointer - // into a vector that may reallocate on subsequent push_back) - vk::Semaphore back_wait_sem = VK_NULL_HANDLE; - if (slot.back_edge_ready) { - back_wait_sem = ggml_vk_create_binary_semaphore(src_ctx)->s; - slot.back_edge_ready = false; - } - vk::Semaphore fwd_sem = ggml_vk_create_binary_semaphore(dst_ctx)->s; - vk::Semaphore back_signal_sem = ggml_vk_create_binary_semaphore(dst_ctx, - vk::ExternalSemaphoreHandleTypeFlagBits::eSyncFd)->s; + size_t aligned_size = ggml_vk_align_size(size, D2D_BUMP_ALIGN); - // Back edge: export previous hop2's signal into the src-side wait semaphore - if (back_wait_sem) { - if (!ggml_vk_d2d_syncfd_export_import(dst->device, slot.last_back_sem, src->device, back_wait_sem)) { - return false; - } + // Periodic hop1 command pool cleanup + if (path.hop1_cmd_pool.buffers_in_use() >= 8) { + vkDeviceWaitIdle(src->device->device); + ggml_vk_command_pool_cleanup(src->device, path.hop1_cmd_pool); } + // Overflow: flush all deferred hop2s, reset bump allocator + if (path.alloc_offset + aligned_size > path.size) { + path.high_water_mark = std::max(path.high_water_mark, path.alloc_offset); + ggml_vk_synchronize(dst_ctx); + vkDeviceWaitIdle(src->device->device); + ggml_vk_command_pool_cleanup(src->device, path.hop1_cmd_pool); + path.alloc_offset = 0; + } + + size_t buf_offset = path.alloc_offset; + path.alloc_offset += aligned_size; + path.high_water_mark = std::max(path.high_water_mark, path.alloc_offset); + + vk::Semaphore fwd_sem = ggml_vk_create_binary_semaphore(dst_ctx)->s; + // Hop 1: src device copies VRAM -> shared buffer { std::lock_guard guard(src->device->mutex); vk_context hop1_ctx = ggml_vk_create_temporary_context(path.hop1_cmd_pool); ggml_vk_ctx_begin(src->device, hop1_ctx); - if (back_wait_sem) { - hop1_ctx->s->wait_semaphores.push_back({ back_wait_sem, 0 }); - } - - VkBufferCopy bc{ src_offset, 0, size }; + VkBufferCopy bc{ src_offset, buf_offset, size }; vkCmdCopyBuffer(hop1_ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)src_side_buf->buffer, 1, &bc); hop1_ctx->s->signal_semaphores.push_back({ path.sem_src, 0 }); @@ -9778,25 +9788,13 @@ static bool ggml_vk_buffer_copy_async_d2d_syncfd( return false; } - // Hop 2: dst device copies shared buffer -> VRAM - { - std::lock_guard guard(dst->device->mutex); - vk_context hop2_ctx = ggml_vk_create_temporary_context(path.hop2_cmd_pool); - ggml_vk_ctx_begin(dst->device, hop2_ctx); + // Hop 2: deferred into dst compute context + vk_context compute_ctx = ggml_vk_get_compute_ctx(dst_ctx); + ggml_vk_ctx_begin(dst->device, compute_ctx); + compute_ctx->s->wait_semaphores.push_back({ fwd_sem, 0 }); - hop2_ctx->s->wait_semaphores.push_back({ fwd_sem, 0 }); - - VkBufferCopy bc2{ 0, dst_offset, size }; - vkCmdCopyBuffer(hop2_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); - - hop2_ctx->s->signal_semaphores.push_back({ back_signal_sem, 0 }); - - ggml_vk_ctx_end(hop2_ctx); - ggml_vk_submit(hop2_ctx, {}); - } - - slot.last_back_sem = back_signal_sem; - slot.back_edge_ready = true; + VkBufferCopy bc2{ buf_offset, dst_offset, size }; + vkCmdCopyBuffer(compute_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); return true; } @@ -9820,7 +9818,7 @@ static bool ggml_vk_buffer_copy_async_d2d( return ggml_vk_buffer_copy_async_d2d_timeline(compute_ctx, dst, dst_offset, src, src_offset, size, path); } - return ggml_vk_buffer_copy_async_d2d_syncfd(src_ctx, dst_ctx, dst, dst_offset, src, src_offset, size, path); + return ggml_vk_buffer_copy_async_d2d_syncfd(dst_ctx, dst, dst_offset, src, src_offset, size, path); } #endif @@ -17088,10 +17086,14 @@ static void ggml_vk_graph_cleanup(ggml_backend_vk_context * ctx) { continue; } if (path.sem_dst_device == ctx->device.get() || path.sem_src_device == ctx->device.get()) { - for (size_t i = 0; i < path.num_slots; i++) { - path.slots[i].last_back_sem = VK_NULL_HANDLE; - path.slots[i].back_edge_ready = false; - } + path.high_water_mark = std::max(path.high_water_mark, path.alloc_offset); + path.alloc_offset = 0; + } + if (path.hop1_device == ctx->device.get() && + path.hop1_cmd_pool.pool && + path.hop1_cmd_pool.buffers_in_use() > 0) { + vkDeviceWaitIdle(ctx->device->device); + ggml_vk_command_pool_cleanup(ctx->device, path.hop1_cmd_pool); } } } From bac136909568b9e071e2b1d59090a8e99db59ff9 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Wed, 10 Jun 2026 15:00:00 +0200 Subject: [PATCH 07/11] swap dmabuf check order --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 37 +++++++++++++++------------- 1 file changed, 20 insertions(+), 17 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index e1bd9faac1..d92fe08b82 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -9387,25 +9387,10 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev bool cross_vendor_nvidia = src_nvidia != dst_nvidia; // 1. dmabuf_p2p — skip if cross-vendor NVIDIA + // Try write direction first: dst exports VRAM, src imports and writes into it. + // PCIe posted writes are generally faster than reads (5-11x on AMD-to-AMD). if (src_dev->external_memory_dma_buf && dst_dev->external_memory_dma_buf && !cross_vendor_nvidia) { - // Try src exports VRAM, dst imports (read direction) vk_buffer exp_buf, imp_buf; - if (ggml_vk_d2d_try_dma_buf(src_dev, dst_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { - if (ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { - path.method = D2D_DMABUF_P2P; - path.reverse_direction = false; - path.slots[0].buf_a = exp_buf; - path.slots[0].buf_b = imp_buf; - path.num_slots = 1; - path.size = VK_D2D_PROBE_SIZE; - GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", - src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); - return path; - } - ggml_vk_destroy_buffer(exp_buf); - ggml_vk_destroy_buffer(imp_buf); - } // Try dst exports VRAM, src imports (write direction) if (ggml_vk_d2d_try_dma_buf(dst_dev, src_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { @@ -9424,6 +9409,24 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev ggml_vk_destroy_buffer(exp_buf); ggml_vk_destroy_buffer(imp_buf); } + + // Try src exports VRAM, dst imports (read direction) + if (ggml_vk_d2d_try_dma_buf(src_dev, dst_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_P2P; + path.reverse_direction = false; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (src exports VRAM)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } } // 2. dmabuf_gtt From bce45fbe90f27ee9c4ad7e2f234ebed19b576ac4 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Thu, 11 Jun 2026 11:37:09 +0200 Subject: [PATCH 08/11] add semi-async staging copy and use it over syncfd for small copies --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 59 ++++++++++++++++++++++++++-- 1 file changed, 56 insertions(+), 3 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index d92fe08b82..0a893feb7e 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -1257,6 +1257,11 @@ struct vk_d2d_path { // reset between evaluations. high_water_mark drives buffer resizing. size_t alloc_offset = 0; size_t high_water_mark = 0; + + // Semi-async: tracks whether a deferred hop2 is pending in the dst + // compute context, so back-to-back copies flush before overwriting + // the shared buffer. + bool semi_async_pending = false; }; #endif @@ -9637,7 +9642,7 @@ static bool ggml_vk_d2d_is_async_capable(vk_device& src_dev, vk_device& dst_dev) std::lock_guard guard(vk_d2d_cache_mutex); auto key = std::make_pair(src_dev.get(), dst_dev.get()); auto it = vk_d2d_cache.find(key); - return it != vk_d2d_cache.end() && it->second.sync_method != D2D_SYNC_NONE; + return it != vk_d2d_cache.end() && it->second.method != D2D_STAGING; } static bool ggml_vk_buffer_copy_async_d2d_timeline( @@ -9802,6 +9807,50 @@ static bool ggml_vk_buffer_copy_async_d2d_syncfd( return true; } +static constexpr size_t D2D_SEMIASYNC_THRESHOLD = 1 * 1024 * 1024; + +static bool ggml_vk_buffer_copy_async_d2d_semiasync( + ggml_backend_vk_context * dst_ctx, + vk_buffer& dst, size_t dst_offset, + vk_buffer& src, size_t src_offset, + size_t size, + vk_d2d_path& path) { + + vk_d2d_path::slot& slot = path.slots[0]; + vk_buffer& src_side_buf = path.reverse_direction ? slot.buf_b : slot.buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? slot.buf_a : slot.buf_b; + + if (path.semi_async_pending) { + ggml_vk_synchronize(dst_ctx); + path.semi_async_pending = false; + } + + // Hop 1: synchronous copy on src device (VRAM -> shared buffer) + { + std::lock_guard guard(src->device->mutex); + vk_context hop1_ctx = ggml_vk_create_temporary_context(src->device->compute_queue.cmd_pool); + ggml_vk_ctx_begin(src->device, hop1_ctx); + + VkBufferCopy bc{ src_offset, 0, size }; + vkCmdCopyBuffer(hop1_ctx->s->buffer->buf, (VkBuffer)src->buffer, (VkBuffer)src_side_buf->buffer, 1, &bc); + + ggml_vk_ctx_end(hop1_ctx); + ggml_vk_submit(hop1_ctx, src->device->fence); + VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "d2d_semiasync hop1 waitForFences"); + src->device->device.resetFences({ src->device->fence }); + } + + // Hop 2: deferred into dst compute context (shared buffer -> dst VRAM) + vk_context compute_ctx = ggml_vk_get_compute_ctx(dst_ctx); + ggml_vk_ctx_begin(dst->device, compute_ctx); + + VkBufferCopy bc2{ 0, dst_offset, size }; + vkCmdCopyBuffer(compute_ctx->s->buffer->buf, (VkBuffer)dst_side_buf->buffer, (VkBuffer)dst->buffer, 1, &bc2); + + path.semi_async_pending = true; + return true; +} + static bool ggml_vk_buffer_copy_async_d2d( ggml_backend_vk_context * src_ctx, ggml_backend_vk_context * dst_ctx, @@ -9812,7 +9861,7 @@ static bool ggml_vk_buffer_copy_async_d2d( vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); - if (path.sync_method == D2D_SYNC_NONE || path.method == D2D_STAGING) { + if (path.method == D2D_STAGING) { return false; } @@ -9821,7 +9870,11 @@ static bool ggml_vk_buffer_copy_async_d2d( return ggml_vk_buffer_copy_async_d2d_timeline(compute_ctx, dst, dst_offset, src, src_offset, size, path); } - return ggml_vk_buffer_copy_async_d2d_syncfd(dst_ctx, dst, dst_offset, src, src_offset, size, path); + if (path.sync_method == D2D_SYNC_SYNCFD && size >= D2D_SEMIASYNC_THRESHOLD) { + return ggml_vk_buffer_copy_async_d2d_syncfd(dst_ctx, dst, dst_offset, src, src_offset, size, path); + } + + return ggml_vk_buffer_copy_async_d2d_semiasync(dst_ctx, dst, dst_offset, src, src_offset, size, path); } #endif From 4d879b9cfa04709307f97be8e15a256fdb40ffc4 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Thu, 18 Jun 2026 11:55:10 +0200 Subject: [PATCH 09/11] use memcpy for small copies across host-visible memory --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 46 +++++++++++++++++++++++----- 1 file changed, 38 insertions(+), 8 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index 0a893feb7e..b353b0ef2a 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -9809,6 +9809,16 @@ static bool ggml_vk_buffer_copy_async_d2d_syncfd( static constexpr size_t D2D_SEMIASYNC_THRESHOLD = 1 * 1024 * 1024; +// Tier 1: both src and dst have rebar — direct memcpy, 0 GPU submits. +// CPU rebar reads are ~1-3 GB/s; below this threshold the memcpy is faster +// than a GPU command buffer submission cycle (~50-100 us). +static constexpr size_t VK_D2D_DIRECT_MEMCPY_THRESHOLD = 128 * 1024; + +// Tier 2: dst has rebar — staging path (1 GPU submit + memcpy via BAR). +// Below this threshold, staging + memcpy-write beats dmabuf's 2 GPU submits +// for rebar destinations. +static constexpr size_t VK_D2D_STAGING_THRESHOLD = 4 * 1024 * 1024; + static bool ggml_vk_buffer_copy_async_d2d_semiasync( ggml_backend_vk_context * dst_ctx, vk_buffer& dst, size_t dst_offset, @@ -9865,6 +9875,13 @@ static bool ggml_vk_buffer_copy_async_d2d( return false; } + // For small copies to rebar destinations, fall back to sync path + // which uses the staging + memcpy-via-BAR optimization + bool dst_mapped = dst->ptr && (dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible); + if (dst_mapped && size <= VK_D2D_STAGING_THRESHOLD) { + return false; + } + if (path.sync_method == D2D_SYNC_TIMELINE) { vk_context compute_ctx = ggml_vk_get_compute_ctx(dst_ctx); return ggml_vk_buffer_copy_async_d2d_timeline(compute_ctx, dst, dst_offset, src, src_offset, size, path); @@ -9903,21 +9920,34 @@ static void ggml_vk_buffer_copy(vk_buffer& dst, size_t dst_offset, vk_buffer& sr ggml_vk_queue_command_pools_cleanup(src->device); } else { VK_LOG_DEBUG("ggml_vk_buffer_copy(MULTI_DEVICE, " << size << ")"); + + bool src_mapped = src->ptr && (src->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible); + bool dst_mapped = dst->ptr && (dst->memory_property_flags & vk::MemoryPropertyFlagBits::eHostVisible); + + // Tier 1: both src and dst have rebar — direct memcpy, no GPU submits + if (src_mapped && dst_mapped && size <= VK_D2D_DIRECT_MEMCPY_THRESHOLD) { + memcpy((uint8_t *)dst->ptr + dst_offset, (const uint8_t *)src->ptr + src_offset, size); + return; + } + #ifdef __linux__ vk_d2d_path& path = ggml_vk_get_d2d_path(src->device, dst->device, size); if (path.method != D2D_STAGING) { - // buf_a is on the src-side device, buf_b is on the dst-side device - // For reverse_direction (dst exports), buf_a is on dst_dev and buf_b is on src_dev - vk_buffer& src_side_buf = path.reverse_direction ? path.slots[0].buf_b : path.slots[0].buf_a; - vk_buffer& dst_side_buf = path.reverse_direction ? path.slots[0].buf_a : path.slots[0].buf_b; + // Tier 2: dst rebar, small/medium — fall through to staging path + // (1 GPU submit + memcpy via BAR is faster than 2 GPU submits) + if (!(dst_mapped && size <= VK_D2D_STAGING_THRESHOLD)) { + // Tier 3: dmabuf path — best bandwidth for large transfers + vk_buffer& src_side_buf = path.reverse_direction ? path.slots[0].buf_b : path.slots[0].buf_a; + vk_buffer& dst_side_buf = path.reverse_direction ? path.slots[0].buf_a : path.slots[0].buf_b; - ggml_vk_buffer_copy(src_side_buf, 0, src, src_offset, size); - ggml_vk_buffer_copy(dst, dst_offset, dst_side_buf, 0, size); - return; + ggml_vk_buffer_copy(src_side_buf, 0, src, src_offset, size); + ggml_vk_buffer_copy(dst, dst_offset, dst_side_buf, 0, size); + return; + } } #endif - // Fallback: staging with CPU memcpy + // Staging fallback: GPU copy to staging + memcpy (or GPU copy) to dst ggml_vk_ensure_sync_staging_buffer(src->device, size); ggml_vk_buffer_copy(src->device->sync_staging, 0, src, src_offset, size); ggml_vk_buffer_write(dst, dst_offset, src->device->sync_staging->ptr, size); From 50e0a2dd8198592ca9b1473385b4b1fa6b923008 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Tue, 25 Aug 2026 11:48:52 +0200 Subject: [PATCH 10/11] fix compile issues --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 12 ++++++------ 1 file changed, 6 insertions(+), 6 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index b353b0ef2a..f8061d000e 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -4186,7 +4186,7 @@ static bool ggml_vk_d2d_try_timeline_sync(vk_device& src_dev, vk_device& dst_dev } try { - path.hop1_cmd_pool.init(src_dev, &src_dev->transfer_queue); + path.hop1_cmd_pool.init(src_dev, src_dev->transfer_queue.get()); } catch (const vk::SystemError& e) { VK_LOG_DEBUG("ggml_vk_d2d_try_timeline_sync: cmd pool/fence creation failed: " << e.what()); if (path.hop1_cmd_pool.pool) { @@ -4239,7 +4239,7 @@ static bool ggml_vk_d2d_try_syncfd_sync(vk_device& src_dev, vk_device& dst_dev, } try { - path.hop1_cmd_pool.init(src_dev, &src_dev->compute_queue); + path.hop1_cmd_pool.init(src_dev, src_dev->compute_queue.get()); } catch (const vk::SystemError& e) { VK_LOG_DEBUG("ggml_vk_d2d_try_syncfd_sync: cmd pool creation failed: " << e.what()); if (path.hop1_cmd_pool.pool) { @@ -9359,14 +9359,14 @@ static bool ggml_vk_d2d_test_copy(vk_device& device, vk_buffer& shared_buf, size } std::lock_guard guard(device->mutex); - vk_context subctx = ggml_vk_create_temporary_context(device->transfer_queue.cmd_pool); + vk_context subctx = ggml_vk_create_temporary_context(device->transfer_queue->cmd_pool); ggml_vk_ctx_begin(device, subctx); VkBufferCopy bc{ 0, 0, size }; vkCmdCopyBuffer(subctx->s->buffer->buf, (VkBuffer)shared_buf->buffer, (VkBuffer)tmp->buffer, 1, &bc); ggml_vk_ctx_end(subctx); try { ggml_vk_submit(subctx, device->fence); - VK_CHECK(device->device.waitForFences({ device->fence }, true, UINT64_MAX), "d2d test copy waitForFences"); + VK_CHECK(device->device.waitForFences({ device->fence }, true, UINT64_MAX), "d2d test copy waitForFences", device); device->device.resetFences({ device->fence }); ggml_vk_queue_command_pools_cleanup(device); ggml_vk_destroy_buffer(tmp); @@ -9838,7 +9838,7 @@ static bool ggml_vk_buffer_copy_async_d2d_semiasync( // Hop 1: synchronous copy on src device (VRAM -> shared buffer) { std::lock_guard guard(src->device->mutex); - vk_context hop1_ctx = ggml_vk_create_temporary_context(src->device->compute_queue.cmd_pool); + vk_context hop1_ctx = ggml_vk_create_temporary_context(src->device->compute_queue->cmd_pool); ggml_vk_ctx_begin(src->device, hop1_ctx); VkBufferCopy bc{ src_offset, 0, size }; @@ -9846,7 +9846,7 @@ static bool ggml_vk_buffer_copy_async_d2d_semiasync( ggml_vk_ctx_end(hop1_ctx); ggml_vk_submit(hop1_ctx, src->device->fence); - VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "d2d_semiasync hop1 waitForFences"); + VK_CHECK(src->device->device.waitForFences({ src->device->fence }, true, UINT64_MAX), "d2d_semiasync hop1 waitForFences", src->device); src->device->device.resetFences({ src->device->fence }); } From e46b3ae5ffd5efb61f56c38215a27b4cb1ca1648 Mon Sep 17 00:00:00 2001 From: Ruben Ortlam Date: Tue, 25 Aug 2026 12:19:26 +0200 Subject: [PATCH 11/11] reverse check order --- ggml/src/ggml-vulkan/ggml-vulkan.cpp | 40 ++++++++++++++-------------- 1 file changed, 20 insertions(+), 20 deletions(-) diff --git a/ggml/src/ggml-vulkan/ggml-vulkan.cpp b/ggml/src/ggml-vulkan/ggml-vulkan.cpp index f8061d000e..6d247599b7 100644 --- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp +++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp @@ -9392,29 +9392,11 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev bool cross_vendor_nvidia = src_nvidia != dst_nvidia; // 1. dmabuf_p2p — skip if cross-vendor NVIDIA - // Try write direction first: dst exports VRAM, src imports and writes into it. - // PCIe posted writes are generally faster than reads (5-11x on AMD-to-AMD). + // Try read direction first: src exports VRAM, dst imports and reads from it. + // DMA reads from peer VRAM tend to be faster than writes on AMD GPUs. if (src_dev->external_memory_dma_buf && dst_dev->external_memory_dma_buf && !cross_vendor_nvidia) { vk_buffer exp_buf, imp_buf; - // Try dst exports VRAM, src imports (write direction) - if (ggml_vk_d2d_try_dma_buf(dst_dev, src_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { - if (ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { - path.method = D2D_DMABUF_P2P; - path.reverse_direction = true; - path.slots[0].buf_a = exp_buf; - path.slots[0].buf_b = imp_buf; - path.num_slots = 1; - path.size = VK_D2D_PROBE_SIZE; - GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", - src_dev->name.c_str(), dst_dev->name.c_str()); - ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); - return path; - } - ggml_vk_destroy_buffer(exp_buf); - ggml_vk_destroy_buffer(imp_buf); - } - // Try src exports VRAM, dst imports (read direction) if (ggml_vk_d2d_try_dma_buf(src_dev, dst_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { if (ggml_vk_d2d_test_copy(dst_dev, imp_buf, VK_D2D_PROBE_SIZE)) { @@ -9432,6 +9414,24 @@ static vk_d2d_path ggml_vk_probe_d2d_path(vk_device& src_dev, vk_device& dst_dev ggml_vk_destroy_buffer(exp_buf); ggml_vk_destroy_buffer(imp_buf); } + + // Try dst exports VRAM, src imports (write direction) + if (ggml_vk_d2d_try_dma_buf(dst_dev, src_dev, VK_D2D_PROBE_SIZE, true, exp_buf, imp_buf)) { + if (ggml_vk_d2d_test_copy(src_dev, imp_buf, VK_D2D_PROBE_SIZE)) { + path.method = D2D_DMABUF_P2P; + path.reverse_direction = true; + path.slots[0].buf_a = exp_buf; + path.slots[0].buf_b = imp_buf; + path.num_slots = 1; + path.size = VK_D2D_PROBE_SIZE; + GGML_LOG_DEBUG("ggml_vulkan: d2d %s -> %s: dmabuf_p2p (dst exports VRAM)\n", + src_dev->name.c_str(), dst_dev->name.c_str()); + ggml_vk_d2d_setup_sync(src_dev, dst_dev, path); + return path; + } + ggml_vk_destroy_buffer(exp_buf); + ggml_vk_destroy_buffer(imp_buf); + } } // 2. dmabuf_gtt