MacKwnoX/NX2navi
1
0
Fork 0
mirror of https://github.com/tangerineturtle1/NX2navi.git synced 2026-11-15 19:19:58 -04:00
Code Issues Packages Projects Releases Wiki Activity

Order the write-back merge's writes after earlier writes of the imported buffers it writes.

Browse Source 
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 <noreply@anthropic.com>
This commit is contained in:
chrg6
2026-11-13 12:32:14 -05:00
parent edd5d04e8c
commit bac47d2edb
1 changed files with 5 additions and 2 deletions
Download Patch File Download Diff File Expand all files Collapse all files

7
runtime/vk_memory.c
View File

@@ -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