From bac47d2edbf1576c3139cefd3b7c745c8e761ea1 Mon Sep 17 00:00:00 2001 From: chrg6 Date: Fri, 13 Nov 2026 12:32:14 -0500 Subject: [PATCH] Order the write-back merge's writes after earlier writes of the imported buffers it writes. The barrier before the merge (vk_memory.c vk_mem_mirror_end) made earlier writes visible to the merge's shader reads only; its shader writes to the imported buffers follow writes there too, and a write after a write needs the write access in the dependency's second scope. On for every device: the barrier was there already, it only covers one access more. Co-Authored-By: Claude Opus 5.5 --- runtime/vk_memory.c | 7 +++++-- 1 file changed, 5 insertions(+), 2 deletions(-) diff --git a/runtime/vk_memory.c b/runtime/vk_memory.c index e8875d95..723b5a3e 100644 --- a/runtime/vk_memory.c +++ b/runtime/vk_memory.c @@ -4322,9 +4322,12 @@ int vk_mem_mirror_end(VkCommandBuffer cb, uint64_t fence) { int rec = 0; lazy_part_check(); // before the merges are recorded: what the last part skipped and has to write back after all if (nsnl) { - // the merge reads the mirror the submission's shaders wrote and writes the imported buffers + // the merge reads the mirror the submission's shaders wrote and writes the imported buffers, which earlier + // commands may have written too: its writes are ordered after theirs as well (a write after a write needs the + // write access on both sides of the dependency) mbarrier(cb, VK_PIPELINE_STAGE_2_ALL_COMMANDS_BIT, VK_ACCESS_2_MEMORY_WRITE_BIT, - VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, VK_ACCESS_2_SHADER_STORAGE_READ_BIT); + VK_PIPELINE_STAGE_2_COMPUTE_SHADER_BIT, + VK_ACCESS_2_SHADER_STORAGE_READ_BIT | VK_ACCESS_2_SHADER_STORAGE_WRITE_BIT); vk_gpu_gap_note(VK_GAP_WB_MERGE, 0); int m = wb_merge_record(cb); if (!m) { // no range list: copy them plainly