From 162bbd77c82dbf6012936cd3e37e9045eea43c80 Mon Sep 17 00:00:00 2001 From: ayo <93222141+countervolts@users.noreply.github.com> Date: Tue, 13 Oct 2026 22:30:41 -0400 Subject: [PATCH] shim: adaptive marker polling, opt-in GPU-side input wait/completion and RGBA8 output - [Interop] MarkerPollUs = auto is the new default: the input-marker wait learns when each feature's inputs become ready, sleeps in <=200 us chunks, yields near the prediction and backs off when late (d4r_input_wait.h); a number keeps fixed pacing (0 yields). d4r_config.py no longer exports 200 when the key is unset - D4R_SHIM_GPU_INPUT_WAIT=1 (experimental, split VRAM, age 0): NGX's null stream waits on a coherent Vulkan word imported into HIP (d4rStreamWaitValue32Eq) instead of CPU polling. The original readback marker stays the authority; a CPU watchdog releases discarded lists and timeouts. Denoiser, verify, capture and dilation frames keep the CPU wait - D4R_SHIM_GPU_COMPLETION=1 (needs GPU_INPUT_WAIT): the bridge writes a completion word right after the native output kernel (d4rNgxWriteValue32), and the game's queue polls it with a bounded compute wait (kernels/support/vulkan_wait.comp), replacing the split list's CPU timeline signal. frame_retired writes the same word on every failure path - D4R_SHIM_NATIVE_COMPLETION=1: a null-stream callback in the bridge signals the split timeline (d4rNgxTimeline*); D4rFrameCompletion::can_signal only releases the oldest outstanding frame - RGBA8 output: tex_common.h adds an 'R4' redirect tag, so K's native output kernel stores packed RGBA8 directly (D4R_SHIM_OUTPUT_RGBA8=1, checked against the kernel manifest by d4rOutputKernelRgba8). Host packing handles RGBA8 too. patches/vkd3d-proton/0004 (VKD3D_D4R_LINEAR_RGBA8=WxH) makes such outputs linear and exportable so the result can be written in place - D4R_SHIM_FUSE_RGBA8_INPUT=1: one Vulkan pass converts and copies RGBA8 color input (vulkan_read.comp mode 5 matches the blit's toward-zero halves, plus a row stride) - profiling: thread CPU time from the Unix bridge (Wine's GetThreadTimes ticks at 10 ms); game-queue D3D12 timestamps (vk_input/vk_dlss_wait/vk_output/vk_total); gpu_eval split into input_gpu_wait and dlss_gpu; D4R_PROFILE_GPU_EVENTS, _TIMESTAMPS, _LOG and _CPU_TIME=0 switch off individual parts for bisection - bridge: D4R_NGX_STREAM_PRIORITY=high|normal runs NGX on a priority stream (off by default; no gain measured) - release_feature: a feature already marked retiring by a failed drain can still be released normally - harness: RGBA8 with motion and jitter scenes (re-quantized each frame), D4R_HARNESS_RGBA8_RAMP, D4R_HARNESS_THRASH_MB to evict L2/Infinity Cache before each evaluation - kernels/native: in_hoist.py and split_output_paths.py (opt-in scheduling experiments); scripts/build_k_ldr_output.sh builds K's LDR output with an optional D4R_LDR_OUTPUT_MAX_VGPR cap - tests: test_input_wait.py and test_latency_optimizations.py; MarkerPollUs and frame-completion tests updated --- config/d4r.ini.default | 7 +- kernels/native/in_hoist.py | 118 +++ kernels/native/split_output_paths.py | 71 ++ kernels/support/vulkan_read.comp | 16 +- kernels/support/vulkan_wait.comp | 12 + kernels/tex/tex_common.h | 22 +- packaging/d4r.ini | 6 +- .../0004-linear-rgba8-output.patch | 35 + scripts/build_k_ldr_output.sh | 41 + scripts/build_vulkan_staging.sh | 1 + scripts/d4r_config.py | 2 +- tests/test_frame_completion.py | 2 + tests/test_input_wait.py | 188 ++++ tests/test_latency_optimizations.py | 204 +++++ tests/test_shim_config.py | 7 +- tools/d3d12_dlss_harness.cpp | 37 +- tools/d4r_frame_completion.h | 5 + tools/d4r_input_wait.h | 30 + tools/d4r_nvngx_shim.cpp | 819 ++++++++++++++++-- tools/d4r_vulkan_staging.h | 53 +- tools/wine_nvcuda_bridge.c | 408 ++++++++- tools/wine_nvcuda_bridge.spec | 14 + 22 files changed, 1989 insertions(+), 109 deletions(-) create mode 100644 kernels/native/in_hoist.py create mode 100644 kernels/native/split_output_paths.py create mode 100644 kernels/support/vulkan_wait.comp create mode 100644 patches/vkd3d-proton/0004-linear-rgba8-output.patch create mode 100755 scripts/build_k_ldr_output.sh create mode 100644 tests/test_input_wait.py create mode 100644 tests/test_latency_optimizations.py create mode 100644 tools/d4r_input_wait.h diff --git a/config/d4r.ini.default b/config/d4r.ini.default index 63dec6e..578c5ac 100644 --- a/config/d4r.ini.default +++ b/config/d4r.ini.default @@ -97,18 +97,19 @@ VramInterop = true InputSync = true ; Faster hand-offs between the game and DLSS (same frame, same result). Need FrameAge = 0 and VRAM ; interop, and a runtime built from this checkout (scripts/install_d4r_runtime.sh). -; Input readiness is checked with CPU sleeps; recorded-but-discarded lists never queue GPU waits. +; Input readiness uses adaptive CPU polling by default. D4R_SHIM_GPU_INPUT_WAIT=1 enables experimental +; null-stream GPU waits for same-frame Super Sampling; diagnostics and host input reads retain CPU waits. ; LinearInputs: DLSS reads the game's inputs in place once the input marker is ready ; DirectOutput: the result is written straight into the game-side buffer (preset K, and only while ; NGX runs the native output kernel; drops NVIDIA's "DLSS SDK" watermark) ; ElideNgxSync: removes NGX's internal CPU waits between its GPU stages ; EvalSync = false: no CPU wait between DLSS and the output copy -; MarkerPollUs: how often (microseconds) the CPU checks for the game's inputs +; MarkerPollUs: auto learns readiness and yields near it; N fixes the poll interval in microseconds LinearInputs = true DirectOutput = true ElideNgxSync = true EvalSync = false -MarkerPollUs = 200 +MarkerPollUs = auto [Engine] ; Experimental. Runs DLSS 4 and 4.5 (models K, L and M) on d4r's own Vulkan inference engine instead of NVIDIA's diff --git a/kernels/native/in_hoist.py b/kernels/native/in_hoist.py new file mode 100644 index 0000000..3e8f99a --- /dev/null +++ b/kernels/native/in_hoist.py @@ -0,0 +1,118 @@ +#!/usr/bin/env python3 +"""Issue independent texture fetches of the translated hiluma input kernel earlier (opt-in experiment). + +usage: in_hoist.py KERNEL.hip ANCHOR_REG HANDLE_REG [HANDLE_REG ...] > OUT.hip + +The input kernel (ptx2hip --ifconv --sink + in_lds2reg.py) is one straight-line body. Its 3x3 color neighbourhood +fetches depend only on the thread's pixel position and kernel parameters, yet the compiler issues them after the +depth -> motion -> history chain, adding serial memory round trips to a latency-bound kernel. This pass moves every +fetch block `{ ::f4 t_ = tex_fetch(rdX, ...); ... }` whose handle register is one of HANDLE_REG's copies, together +with the backward slice of statements computing its operands, to right after the statement that assigns +ANCHOR_REG, followed by a scheduling barrier (instruction order only). Nothing else moves and no statement +changes, so every value is computed exactly as before. + +Refuses unless: every register the moved statements assign is assigned exactly once in the kernel; every register +they read is assigned before the anchor (or by another moved statement, earlier in order); there is no memory store +and no control flow (label/goto/return) between the anchor and the last moved statement. +""" +import re +import sys + +REG = re.compile(r'\b(?:rd|rs|r|f|p|shv\d+_|dpv\d+_)\d+\b|\bl\d+_\w+\b') +ASSIGN = re.compile(r'^\s*([A-Za-z_]\w*)\s*=(?!=)') + + +def stmts(body_lines): + """group lines into statements: fetch blocks '{ ... }' become one statement""" + out, i = [], 0 + while i < len(body_lines): + line = body_lines[i] + if line.strip().startswith('{ ::f4 t_ ='): + j = i + while body_lines[j].strip() != '}': + j += 1 + out.append(body_lines[i:j + 1]) + i = j + 1 + else: + out.append([line]) + i += 1 + return out + + +def defs_uses(st): + text = ''.join(st) + if text.strip().startswith('{ ::f4 t_ ='): + defs = set(re.findall(r'^\s*([A-Za-z_]\w*) = t_\[\d\];', text, re.M)) + call = text.split(';', 1)[0].replace('::f4', '') + return defs, set(REG.findall(call)) + m = ASSIGN.match(text) + if m: + lhs = m.group(1) + rhs = text[m.end():] + return {lhs}, set(REG.findall(rhs)) + return set(), set(REG.findall(text)) + + +def main(): + src = open(sys.argv[1]).read().split('\n') + anchor, handles = sys.argv[2], set(sys.argv[3:]) + start = next(i for i, l in enumerate(src) if l.startswith('extern "C" __attribute__((global))')) + 2 + end = max(i for i, l in enumerate(src) if l.strip() == '}') + S = stmts([l + '\n' for l in src[start:end]]) + info = [defs_uses(s) for s in S] + assigned = {} + for k, (d, _) in enumerate(info): + for r in d: + assigned.setdefault(r, []).append(k) + # handle copies: rdY = (uint64_t)(rdX) chains from the given handle registers + alias = set(handles) + for k, s in enumerate(S): + m = re.match(r'^\s*(rd\d+) = \(uint64_t\)\((rd\d+)\);$', s[0]) + if m and m.group(2) in alias: + alias.add(m.group(1)) + ia = [k for k, (d, _) in enumerate(info) if anchor in d] + if len(ia) != 1: + raise SystemExit(f'anchor {anchor} assigned {len(ia)} times') + ia = ia[0] + fetches = [k for k, s in enumerate(S) if s[0].strip().startswith('{ ::f4 t_ = tex_fetch(') + and re.search(r'tex_fetch\((rd\d+)', s[0]).group(1) in alias and k > ia] + if not fetches: + raise SystemExit('no fetch to move') + need, move = set(), set(fetches) + for k in fetches: + need |= info[k][1] + while need: + r = need.pop() + ks = assigned.get(r, []) + if not ks: + continue # parameters / builtins + if len(ks) != 1: + raise SystemExit(f'register {r} is assigned {len(ks)} times') + k = ks[0] + if k <= ia or k in move: + continue + move.add(k) + need |= info[k][1] + last = max(move) + for k in range(ia + 1, last + 1): + text = ''.join(S[k]) + if re.search(r'^\s*\*\(|\bgoto\b|\breturn\b|^L__BB', text, re.M): + raise SystemExit(f'store or control flow between anchor and moved code: {text.strip()[:80]}') + for k in sorted(move): + for r in info[k][0]: + if len(assigned[r]) != 1: + raise SystemExit(f'moved register {r} assigned more than once') + for r in info[k][1]: + ks = assigned.get(r, []) + if ks and not (ks[0] <= ia or ks[0] in move and ks[0] < k): + raise SystemExit(f'moved statement reads {r} defined later') + # a scheduling barrier keeps the compiler from sinking the fetches back to their uses (order only) + barrier = [' __builtin_amdgcn_sched_barrier(0); // in_hoist: fetches above are issued here\n'] + body = ''.join(''.join(S[k]) for k in range(0, ia + 1)) + ''.join(''.join(S[k]) for k in sorted(move)) + \ + ''.join(barrier) + ''.join(''.join(S[k]) for k in range(ia + 1, len(S)) if k not in move) + sys.stderr.write(f'in_hoist: moved {len(move)} statements ({len(fetches)} fetches) after {anchor}\n') + sys.stdout.write('\n'.join(src[:start]) + '\n' + body + '\n'.join(src[end:])) + + +if __name__ == '__main__': + main() diff --git a/kernels/native/split_output_paths.py b/kernels/native/split_output_paths.py new file mode 100644 index 0000000..4812377 --- /dev/null +++ b/kernels/native/split_output_paths.py @@ -0,0 +1,71 @@ +#!/usr/bin/env python3 +"""Split a translated hiluma output kernel into its two per-block paths (opt-in experiment). + +usage: split_output_paths.py KERNEL.hip > KERNEL_split.hip + +NVIDIA's hiluma_engine_output_* kernels pick one of two code paths per block: path B (an LDS tile of the +block's render footprint) when the footprint is under 24 texels in both axes, which holds for every block at +Quality .. Ultra Performance, else path A (direct fetches, DLAA). The choice depends only on the block index and +kernel parameters, so it is uniform per block, and blocks share nothing. Path A needs more registers than path B +(LDR mvlo: 101 vs 88 VGPRs), and one kernel carries the larger allocation for every block. + +This pass emits the same translated body twice: + NAME runs path B and returns where path A would start; + NAME_post1 runs path A and returns where path B would start. +The ZLUDA/MicroCUDA native hook launches NAME_post1 right after NAME on the same stream with the same parameters +and grid (d4r_post1_grid_x/_y = 0), with the 16x16 blocks the PTX hard-codes (tid.y << 4). Each block executes +exactly the instructions it executed before; only the block order across the two paths changes, which no output +depends on (blocks never read what other blocks write in the same launch). + +The pass refuses anything but the expected entry structure: + pA = ((int32_t)rX < (int32_t)24); pB = ((int32_t)rY < (int32_t)24); pC = (pA && pB); + if (pC) goto LB; + goto LA; +""" +import re +import sys + + +def split(source): + head, sep, main = source.partition('extern "C" __attribute__((global))') + if not sep or 'extern "C" __attribute__((global))' in main: + raise SystemExit('expected exactly one kernel') + m = re.match(r'\s*__attribute__\(\(amdgpu_flat_work_group_size\(1, 256\)\)\) void (\w+)\((\w+)\)', main) + if not m: + raise SystemExit('unexpected kernel signature') + name = m.group(1) + pattern = re.compile(r'^ (p\d+) = \(\(int32_t\)r\d+ < \(int32_t\)24\);\n' + r' (p\d+) = \(\(int32_t\)r\d+ < \(int32_t\)24\);\n' + r' (p\d+) = \((p\d+) && (p\d+)\);\n' + r' if \((p\d+)\) goto (L__BB0_\d+);\n' + r' goto (L__BB0_\d+);\n', re.M) + found = list(pattern.finditer(main)) + if len(found) != 1: + raise SystemExit(f'path selection found {len(found)} times, expected once') + g = found[0] + pa, pb, pc, x, y, pc2, path_b, path_a = g.groups() + if {x, y} != {pa, pb} or pc != pc2: + raise SystemExit('path selection predicate has an unexpected form') + # the predicate must not be reassigned anywhere (it is the block's path choice) + if len(re.findall(rf'^ {pc} = ', main, re.M)) != 1: + raise SystemExit(f'{pc} is assigned more than once') + sel = g.group(0) + lines = sel.splitlines(keepends=True) + main_b = main.replace(sel, ''.join(lines[:4]) + ' return; // path A blocks: NAME_post1\n', 1) + main_a = main.replace(sel, ''.join(lines[:3]) + f' if ({pc}) return; // path B blocks: NAME\n' + + f' goto {path_a};\n', 1) + main_a = main_a.replace(f' void {name}(', f' void {name}_post1(', 1) + globals_ = ('\n// d4r native hook: launch NAME_post1 after NAME with the launch grid and the PTX\'s 16x16 blocks\n' + 'extern "C" __attribute__((device, used)) uint32_t d4r_post1_grid_x = 0;\n' + 'extern "C" __attribute__((device, used)) uint32_t d4r_post1_grid_y = 0;\n' + 'extern "C" __attribute__((device, used)) uint32_t d4r_post1_block_x = 16;\n' + 'extern "C" __attribute__((device, used)) uint32_t d4r_post1_block_y = 16;\n' + 'extern "C" __attribute__((device, used)) uint32_t d4r_post1_block_z = 1;\n') + banner = '// split_output_paths.py: NAME = path B only, NAME_post1 = path A only (per-block choice unchanged)\n' + return (banner + head + globals_ + sep + main_b + '\n' + sep + main_a, path_b, path_a) + + +if __name__ == '__main__': + text, b, a = split(open(sys.argv[1]).read()) + sys.stderr.write(f'split: path B {b}, path A {a}\n') + sys.stdout.write(text) diff --git a/kernels/support/vulkan_read.comp b/kernels/support/vulkan_read.comp index ec29176..98e04e1 100644 --- a/kernels/support/vulkan_read.comp +++ b/kernels/support/vulkan_read.comp @@ -2,7 +2,7 @@ layout(local_size_x=8, local_size_y=8) in; layout(set=0, binding=0) uniform sampler2D sourceImage; layout(std430, set=0, binding=1) writeonly buffer Pixels { uint pixels[]; }; -layout(push_constant) uniform Size { uint width; uint height; uint mode; } size; +layout(push_constant) uniform Size { uint width; uint height; uint mode; uint stride; } size; // Match the CPU canonical conversion, including round-to-even and subnormals. uint halfBits(float value) { uint bits = floatBitsToUint(value); @@ -26,14 +26,24 @@ uint halfBits(float value) { if (remainder > halfway || (remainder == halfway && (m & 1u) != 0u)) ++result; return result; } +// Mode 5 is only for RGBA8 UNORM: values are zero or normal positive halves. Preserve +// the blit's toward-zero conversion; other modes keep the canonical round-to-even path. +uint halfRtz(float value) { + uint b = floatBitsToUint(value); + if (value == 0.0) return 0u; + return (((b >> 23) - 112u) << 10) | ((b & 0x7fffffu) >> 13); +} void main() { uvec2 p = gl_GlobalInvocationID.xy; if (p.x >= size.width || p.y >= size.height) return; ivec2 extent = textureSize(sourceImage, 0); - uint index = p.y * size.width + p.x; + uint index = p.y * size.stride + p.x; vec4 value = all(lessThan(ivec2(p), extent)) ? texelFetch(sourceImage, ivec2(p), 0) : vec4(0.0); - if (size.mode == 1u) { // canonical RGBA16F color + if (size.mode == 5u) { + pixels[index * 2u] = halfRtz(value.x) | (halfRtz(value.y) << 16); + pixels[index * 2u + 1u] = halfRtz(value.z) | (halfRtz(value.w) << 16); + } else if (size.mode == 1u) { // canonical RGBA16F color pixels[index * 2u] = halfBits(value.x) | (halfBits(value.y) << 16); pixels[index * 2u + 1u] = halfBits(value.z) | (halfBits(value.w) << 16); } else if (size.mode == 3u) { // canonical RG16F motion diff --git a/kernels/support/vulkan_wait.comp b/kernels/support/vulkan_wait.comp new file mode 100644 index 0000000..e70c7f4 --- /dev/null +++ b/kernels/support/vulkan_wait.comp @@ -0,0 +1,12 @@ +#version 450 +// D4R_SHIM_GPU_COMPLETION: holds the game's queue until DLSS has written this frame's completion word (HIP +// hipStreamWriteValue32 after the output, or the CPU on every failure path). Bounded: after `limit` polls the +// shader returns regardless, so a lost write can never hang the GPU (the frame then shows its previous output). +layout(local_size_x = 1) in; +layout(std430, set = 0, binding = 0) coherent volatile readonly buffer Words { uint words[]; }; +layout(push_constant) uniform Wait { uint index; uint value; uint limit; } wait; +void main() { + for (uint poll = 0u; poll < wait.limit; ++poll) + if (words[wait.index] == wait.value) + return; +} diff --git a/kernels/tex/tex_common.h b/kernels/tex/tex_common.h index 43a3952..66620c5 100644 --- a/kernels/tex/tex_common.h +++ b/kernels/tex/tex_common.h @@ -371,15 +371,19 @@ extern "C" __attribute__((device)) int __ockl_image_width_2D(tsharp_t* image); extern "C" __attribute__((device)) int __ockl_image_height_2D(tsharp_t* image); __attribute__((device)) static inline uint32_t redirect_pitch(uint64_t surface) { - // 'R2' (RGBA16F texels) or 'R3' (R10G10B10A2_UNORM texels) tag in the high half (other handles may have + // 'R2' (RGBA16F), 'R3' (RGB10A2) or 'R4' (RGBA8) tag in the high half (other handles may have // data there), row pitch / 8 in the low half const uint32_t word = ((const __attribute__((address_space(4))) uint32_t*)surface)[21]; - return (word >> 16) == 0x5232u || (word >> 16) == 0x5233u ? (word & 0xffffu) << 3 : 0u; + return (word >> 16) >= 0x5232u && (word >> 16) <= 0x5234u ? (word & 0xffffu) << 3 : 0u; } __attribute__((device)) static inline bool redirect_rgb10a2(uint64_t surface) { return (((const __attribute__((address_space(4))) uint32_t*)surface)[21] >> 16) == 0x5233u; } +__attribute__((device)) static inline bool redirect_rgba8(uint64_t surface) +{ + return (((const __attribute__((address_space(4))) uint32_t*)surface)[21] >> 16) == 0x5234u; +} // One channel as the game-side path produces it: the f32 value rounded toward zero to a half (the RGBA16F // store), then converted to UNORM by a blit (NaN and negatives 0, clamped at 1, rounded to nearest). __attribute__((device)) static inline uint32_t redirect_unorm(uint16_t half, float scale) @@ -422,6 +426,10 @@ DEV __attribute__((always_inline)) void d4r_sust_p_v4b32(uint64_t surface, int32 *D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + 4 * x) = redirect_unorm((uint16_t)lo, 1023.0f) | redirect_unorm((uint16_t)(lo >> 16), 1023.0f) << 10 | redirect_unorm((uint16_t)hi, 1023.0f) << 20 | redirect_unorm((uint16_t)(hi >> 16), 3.0f) << 30; + else if (redirect_rgba8(surface)) + *D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + 4 * x) = + redirect_unorm((uint16_t)lo, 255.0f) | redirect_unorm((uint16_t)(lo >> 16), 255.0f) << 8 | + redirect_unorm((uint16_t)hi, 255.0f) << 16 | redirect_unorm((uint16_t)(hi >> 16), 255.0f) << 24; else *D4R_REDIRECT_PTR(uint2_t, redirect_row(surface, y) + 8 * x) = (uint2_t){lo, hi}; } @@ -526,9 +534,15 @@ DEV __attribute__((always_inline)) void d4r_sust_b32(uint64_t surface, int32_t x tsharp_t* image = (tsharp_t*)surface; if (redirect_pitch(surface) != 0) { + // Raw RGBA16F stores carry two half channels at each four-byte offset. Packing those channels + // separately keeps the original store ownership (no read-modify-write or cross-thread race). + if (redirect_rgba8(surface) && x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y)) + *D4R_REDIRECT_PTR(uint16_t, redirect_row(surface, y) + (x >> 1)) = + (uint16_t)(redirect_unorm((uint16_t)data, 255.0f) | + redirect_unorm((uint16_t)(data >> 16), 255.0f) << 8); // raw bits at byte offset x (RGBA16F rows are 8 bytes per pixel); the shim only asks for packed - // RGB10A2 texels from kernels without raw stores - if (!redirect_rgb10a2(surface) && x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y)) + // RGB10A2 texels from kernels without raw stores; RGBA8 preserves each half-pair store + if (!redirect_rgb10a2(surface) && !redirect_rgba8(surface) && x >= 0 && (x & 3) == 0 && redirect_inside(surface, x >> 3, y)) *D4R_REDIRECT_PTR(uint32_t, redirect_row(surface, y) + x) = data; return; } diff --git a/packaging/d4r.ini b/packaging/d4r.ini index 3d0d4bb..0acb400 100644 --- a/packaging/d4r.ini +++ b/packaging/d4r.ini @@ -62,10 +62,10 @@ NativeFp8 = on ; Requires render- or display-resolution, unjittered motion vectors and same-frame linear VRAM inputs. Restart to change. ; Unsupported features keep original vectors; larger radii may affect disocclusion edges. MotionVectorDilation = 0 -; Keep inputs and output in video memory. Input and output readiness waits use CPU sleeps. +; Keep inputs and output in video memory. Readiness uses adaptive CPU polling by default. VramInterop = true -; Microseconds between checks for the game's input marker (default 200; 0 yields). -MarkerPollUs = 200 +; auto: adaptive input-marker polling; N: fixed interval in microseconds (0 yields). +MarkerPollUs = auto LinearInputs = true ElideNgxSync = true ; Lets DLSS's last kernel write straight into the game's output buffer (DLSS 4 / K only). It takes diff --git a/patches/vkd3d-proton/0004-linear-rgba8-output.patch b/patches/vkd3d-proton/0004-linear-rgba8-output.patch new file mode 100644 index 0000000..426c415 --- /dev/null +++ b/patches/vkd3d-proton/0004-linear-rgba8-output.patch @@ -0,0 +1,35 @@ +vkd3d-proton (d4r, opt-in): VKD3D_D4R_LINEAR_RGBA8=WxH also creates RGBA8 UNORM textures of exactly that size with +unordered access linear and exportable (with VKD3D_D4R_LINEAR_TEXTURES=1), so the d4r shim can write DLSS's packed +RGBA8 output in place (D4R_SHIM_INPLACE=1) instead of copying it. Every selected texture is logged (warn). +Applies on top of 0001-0003. + +diff --git a/libs/vkd3d/resource.c b/libs/vkd3d/resource.c +--- a/libs/vkd3d/resource.c ++++ b/libs/vkd3d/resource.c +@@ -720,7 +720,25 @@ + desc->DepthOrArraySize != 1 || desc->SampleDesc.Count != 1 || + (desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_DEPTH_STENCIL)) + return false; +- return format->vk_format == VK_FORMAT_R16G16B16A16_SFLOAT || format->vk_format == VK_FORMAT_A2B10G10R10_UNORM_PACK32; ++ if (format->vk_format == VK_FORMAT_R16G16B16A16_SFLOAT || format->vk_format == VK_FORMAT_A2B10G10R10_UNORM_PACK32) ++ return true; ++ /* VKD3D_D4R_LINEAR_RGBA8=WxH (opt-in, with VKD3D_D4R_LINEAR_TEXTURES=1): also RGBA8 UNORM textures of exactly ++ * that size with unordered access, the shape of a DLSS output, so the shim can write DLSS's packed RGBA8 result ++ * in place. Other RGBA8 targets keep optimal tiling. Every selected texture is logged. */ ++ if (format->vk_format == VK_FORMAT_R8G8B8A8_UNORM && (desc->Flags & D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS)) ++ { ++ static unsigned int selected; ++ unsigned int width = 0, height = 0; ++ char size[32]; ++ if (vkd3d_get_env_var("VKD3D_D4R_LINEAR_RGBA8", size, sizeof(size)) && ++ sscanf(size, "%ux%u", &width, &height) == 2 && desc->Width == width && desc->Height == height) ++ { ++ WARN("d4r: linear exportable RGBA8 texture #%u (%ux%u, flags %#x).\n", ++selected, width, height, ++ desc->Flags); ++ return true; ++ } ++ } ++ return false; + } + + static HRESULT vkd3d_get_image_create_info(struct d3d12_device *device, diff --git a/scripts/build_k_ldr_output.sh b/scripts/build_k_ldr_output.sh new file mode 100755 index 0000000..f0cfc2d --- /dev/null +++ b/scripts/build_k_ldr_output.sh @@ -0,0 +1,41 @@ +#!/usr/bin/env bash +# Compile K's LDR output using the same PTX semantics as the HDR native route. The packed-store +# sidecar authenticates the produced code bytes; it is not a substitute for temporal validation. +set -euo pipefail +ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" +OUT="$(realpath -m "${1:?usage: build_k_ldr_output.sh OUT_DIR [KERNEL]}")" +NAME="${2:-hiluma_engine_output_depthinv_mvlo_ldr_max_v2_rel}" +[[ "$NAME" =~ ^hiluma_engine_output_depth(inv|reg)_mv(hi|lo)_ldr_max_v2_rel$ ]] || exit 2 +: "${D4R_DLSS_DLL:?set D4R_DLSS_DLL to the matching NVIDIA DLSS library}" +mkdir -p "$OUT" +WORK="$(mktemp -d)" +trap 'rm -rf "$WORK"' EXIT +python3 "$ROOT/kernels/tools/extract_dlss_ptx.py" "$D4R_DLSS_DLL" "$WORK/ptx" >/dev/null +PTX="$(rg -l "\.entry[[:space:]]+$NAME[[:space:]]*\(" "$WORK/ptx"/*.ptx)" +[[ -n "$PTX" && "$PTX" != *$'\n'* ]] || { echo 'kernel must appear in exactly one module' >&2; exit 2; } +# LDR uses another CFG/parameter layout than HDR: do not apply HDR's tile rewrite by register name. +python3 "$ROOT/kernels/native/ptx2hip.py" --as1 "$PTX" "$NAME" > "$WORK/output.hip" +# D4R_LDR_OUTPUT_MAX_VGPR=N (opt-in): cap the kernel's VGPRs. The DLAA path (path A, 100 VGPRs alone) otherwise sets +# the allocation of every block (101 VGPRs, 7 waves/SIMD); with 96 its spills stay in path A and Quality..UP blocks +# (path B, 88 VGPRs alone) run at 8 waves/SIMD: depthinv_mvlo output 0.825 -> 0.776 ms at 1440p, DLAA 0.840 -> +# 0.812, every frame byte-identical (private/notes/k1440-implementation.md, O1). Lower caps spill in path B. +if [[ -n "${D4R_LDR_OUTPUT_MAX_VGPR:-}" ]]; then + [[ "$D4R_LDR_OUTPUT_MAX_VGPR" =~ ^[0-9]+$ ]] || { echo 'D4R_LDR_OUTPUT_MAX_VGPR must be a number' >&2; exit 2; } + sed -i "s/__attribute__((amdgpu_flat_work_group_size(1, 256))) void $NAME(/__attribute__((amdgpu_flat_work_group_size(1, 256), amdgpu_num_vgpr($D4R_LDR_OUTPUT_MAX_VGPR))) void $NAME(/" "$WORK/output.hip" + grep -q "amdgpu_num_vgpr($D4R_LDR_OUTPUT_MAX_VGPR)" "$WORK/output.hip" || { echo 'VGPR cap not applied' >&2; exit 2; } +fi +bash "$ROOT/kernels/native/build_native.sh" "$WORK/output.hip" "$OUT/$NAME.hsaco" \ + -DD4R_NATIVE_F2I -DD4R_AS1_STORES +python3 - "$OUT" "$NAME" <<'PY' +from pathlib import Path +import sys +out, name = Path(sys.argv[1]), sys.argv[2] +value = 14695981039346656037 +for byte in (out / (name + '.hsaco')).read_bytes(): + value = ((value ^ byte) * 1099511628211) & ((1 << 64) - 1) +path = out / 'd4r-rgba8.txt' +rows = path.read_text().splitlines() if path.exists() else ['# RGBA8 store capability: NAME FNV1a64 of code object'] +rows = [row for row in rows if not row.startswith(name + ' ')] +path.write_text('\n'.join(rows + [f'{name} {value:016x}']) + '\n') +PY +python3 "$ROOT/kernels/tools/kernel_manifest.py" "$OUT" "$D4R_DLSS_DLL" diff --git a/scripts/build_vulkan_staging.sh b/scripts/build_vulkan_staging.sh index f536d42..d0496fc 100644 --- a/scripts/build_vulkan_staging.sh +++ b/scripts/build_vulkan_staging.sh @@ -4,6 +4,7 @@ ROOT="$(cd "$(dirname "${BASH_SOURCE[0]}")/.." && pwd)" OUT="${1:-$ROOT/build/vulkan-staging}" mkdir -p "$OUT" glslangValidator -V --target-env vulkan1.0 -o "$OUT/read.spv" "$ROOT/kernels/support/vulkan_read.comp" >/dev/null +glslangValidator -V --target-env vulkan1.0 -o "$OUT/wait.spv" "$ROOT/kernels/support/vulkan_wait.comp" >/dev/null for format in rgba16f rgba32f r11f_g11f_b10f rgb10_a2 rgba8; do glslangValidator -V --target-env vulkan1.0 -DOUTPUT_FORMAT="$format" \ -o "$OUT/write_$format.spv" "$ROOT/kernels/support/vulkan_write.comp" >/dev/null diff --git a/scripts/d4r_config.py b/scripts/d4r_config.py index f5c0028..0d2ee9f 100755 --- a/scripts/d4r_config.py +++ b/scripts/d4r_config.py @@ -161,7 +161,7 @@ def main(argv): value = flag("Interop", key) if value is not None: env[var] = "1" if value else "0" - value = get("Interop", "MarkerPollUs") or "200" + value = get("Interop", "MarkerPollUs") dilation = get("Interop", "MotionVectorDilation") if dilation is not None and dilation not in ("0", "1", "2"): fail("[Interop] MotionVectorDilation must be 0, 1 or 2") diff --git a/tests/test_frame_completion.py b/tests/test_frame_completion.py index a87e3d6..ac140dd 100644 --- a/tests/test_frame_completion.py +++ b/tests/test_frame_completion.py @@ -30,8 +30,10 @@ struct Feature { int splitSemaphore = 1; bool engineNetTimeline = false; // the Vulkan timeline of this test, not the native network's int engineNet = 0; + void* ngxTimeline = nullptr; }; static struct { int (*signalTimeline)(int, uint64_t) = nullptr; } g_engineNet; +static struct { struct { int (*timelineSignal)(void*, uint64_t) = nullptr; } cu; } g; static uint64_t gpuTimeline = 0; static bool failSignal = false; static int signal(int, const VkSemaphoreSignalInfo* info) { diff --git a/tests/test_input_wait.py b/tests/test_input_wait.py new file mode 100644 index 0000000..3ac22cc --- /dev/null +++ b/tests/test_input_wait.py @@ -0,0 +1,188 @@ +"""CPU-only input-wait cancellation, polling and profiling contracts. No HIP or Wine runs.""" +from pathlib import Path +import subprocess +import tempfile +import unittest + +ROOT = Path(__file__).resolve().parents[1] + + +def function(source, signature): + start = source.index(signature) + opening = source.index('{', start) + depth = 1 + end = opening + 1 + while depth: + depth += (source[end] == '{') - (source[end] == '}') + end += 1 + return source[start:end] + + +def run(source): + with tempfile.TemporaryDirectory() as tmp: + runner = Path(tmp) / 'input-wait' + subprocess.run(['c++', '-x', 'c++', '-std=c++20', '-pthread', '-Wall', '-Wextra', '-Werror', + '-I', str(ROOT / 'tools'), '-o', str(runner), '-'], + input=source, text=True, check=True, capture_output=True) + subprocess.run([str(runner)], check=True, timeout=5) + + +class InputWaitTests(unittest.TestCase): + def test_bridge_queues_exact_null_stream_wait_or_reports_unsupported(self): + bridge = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text() + wait = function(bridge, 'CUresult WINAPI d4rStreamWaitValue32Eq(') + run(r''' +#include +#include +#include +#define WINAPI +using CUresult = int; +using CUdeviceptr = uint64_t; +using HIP_STREAM_WAIT_VALUE32_FN = int(*)(void*, void*, uint32_t, unsigned, uint32_t); +constexpr int CUDA_SUCCESS = 0, CUDA_ERROR_NOT_SUPPORTED = 801; +int context_setup_result = 0, supported = 1, waitResult = 0, queued = 0; +bool hasWait = true; +void ensure_context() {} +void* d4r_ngx_stream(void* stream) { return stream; } // D4R_NGX_STREAM_PRIORITY off +void tracef(const char*, ...) {} +bool trace_verbose() { return false; } +int device(int* out) { *out = 3; return 0; } +int attribute(int* out, int attr, int dev) { assert(attr == 10013 && dev == 3); *out = supported; return 0; } +int queue(void* stream, void* ptr, uint32_t value, unsigned flags, uint32_t mask) { + assert(stream == nullptr && ptr == reinterpret_cast(0x10008)); + assert(value == 0xffffffffu && flags == 1 && mask == 0xffffffffu); + if (waitResult == 0) ++queued; + return waitResult; +} +void* hip_symbol(const char* name) { + if (strcmp(name, "hipStreamWaitValue32") == 0) return hasWait ? reinterpret_cast(queue) : nullptr; + if (strcmp(name, "hipDeviceGetAttribute") == 0) return reinterpret_cast(attribute); + if (strcmp(name, "hipGetDevice") == 0) return reinterpret_cast(device); + return nullptr; +} +''' + wait + r''' +int main() { + assert(d4rStreamWaitValue32Eq(0x10008, 0xffffffffu) == 0 && queued == 1); + supported = 0; + assert(d4rStreamWaitValue32Eq(0x10008, 0xffffffffu) == 801 && queued == 1); + supported = 1; hasWait = false; + assert(d4rStreamWaitValue32Eq(0x10008, 0xffffffffu) == 801 && queued == 1); + hasWait = true; waitResult = 1; + assert(d4rStreamWaitValue32Eq(0x10008, 0xffffffffu) == 801 && queued == 1); + context_setup_result = 700; + assert(d4rStreamWaitValue32Eq(0x10008, 0xffffffffu) == 700 && queued == 1); +} +''') + + def test_readiness_discard_retirement_timeout_wrap_and_adaptive_pacing(self): + run(r''' +#include "d4r_input_wait.h" +#include +int main() { + using A = D4rInputWaitAbort; + assert(d4r_input_wait_abort(7, 6, 6, false, 100) == A::None); + assert(d4r_input_wait_abort(7, 6, 8, false, 100) == A::Discarded); + assert(d4r_input_wait_abort(7, 7, 8, false, 100) == A::None); // slot recheck wins + assert(d4r_input_wait_abort(7, 7, 7, true, 100) == A::Retired); + assert(d4r_input_wait_abort(7, 6, 6, false, 5000001) == A::Timeout); + assert(d4r_input_wait_abort(7, 7, 7, false, 5000001) == A::Timeout); // broken wait after input + assert(d4r_input_wait_abort(0xffffffffu, 3, 0, false, 100) == A::Discarded); + assert(d4r_input_wait_abort(0, 0xffffffffu, 0xffffffffu, false, 100) == A::None); + assert(d4r_input_wait_abort(0, 0xffffffffu, 1, false, 100) == A::Discarded); + assert(d4r_input_wait_abort(0, 0, 1, false, 100) == A::None); + assert(d4r_marker_poll_delay_us(2000, 0) == 200); + assert(d4r_marker_poll_delay_us(2000, 1800) == 50); + assert(d4r_marker_poll_delay_us(2000, 1850) == 0); + assert(d4r_marker_poll_delay_us(2000, 2149) == 0); + assert(d4r_marker_poll_delay_us(2000, 2150) == 20); + assert(d4r_marker_poll_delay_us(2000, 3750) == 200); + assert(d4r_marker_poll_delay_us(0, 5000) == 200); // no warm prediction +} +''') + + def test_actual_release_and_drain_helpers_preserve_real_markers_and_ownership(self): + shim = (ROOT / 'tools/d4r_nvngx_shim.cpp').read_text() + release = function(shim, 'static void release_ngx_input_wait_locked(') + complete = function(shim, 'static bool complete_ngx_input_wait(Feature& feature, uint32_t frame, bool abort)\n{') + run(r''' +#include +#include +#include +#include +#include +void logf(const char*, ...) {} +void MemoryBarrier() { std::atomic_thread_fence(std::memory_order_seq_cst); } +struct Feature { + std::atomic ngxGpuWait{true}, retiring{false}; + void* ngxTimeline = nullptr; + std::mutex ngxWaitMutex; + bool ngxWaitActive = true, ngxWaitCancelled = false; + uint32_t ngxWaitFrame = 7; + int ngxWaitSlot = 1; + struct { void* mapped; } ngxInputSync; + unsigned evaluatedFrames = 100; +}; +uint32_t words[4] = {0, 3, 4, 5}, markers[4] = {6, 3, 4, 5}; +Feature* current; +int drains = 0, drainResult = 0; +int sync() { + ++drains; + assert(words[2] == 7); // drain may only run after the CPU released the blocked queue + assert(current->ngxWaitActive); // ticket and input ownership survive the drain + return drainResult; +} +struct { struct { int(*ctxSynchronize)() = sync; void(*timelineCancel)(void*) = nullptr; } cu; } g; +''' + release + '\n' + complete + r''' +int main() { + Feature f; f.ngxInputSync.mapped = words; current = &f; + assert(complete_ngx_input_wait(f, 8, true)); // another frame cannot release this ticket + assert(drains == 0 && words[2] == 4); + // An independent thread can release a blocked host submission without invoking HIP. + auto blocked = std::async(std::launch::async, [&] { + for (;;) { + { std::lock_guard lock(f.ngxWaitMutex); if (words[2] == 7) return; } + std::this_thread::yield(); + } + }); + { std::lock_guard lock(f.ngxWaitMutex); + release_ngx_input_wait_locked(f, "discard", false); } + blocked.get(); + assert(drains == 0 && f.ngxWaitCancelled && f.ngxGpuWait); + assert(words[0] == 0 && words[1] == 3 && words[3] == 5); + assert(markers[0] == 6 && markers[2] == 4); // never counterfeit game readiness + assert(!complete_ngx_input_wait(f, 7, true)); + assert(drains == 1 && !f.ngxWaitActive && f.evaluatedFrames == 0); + assert(complete_ngx_input_wait(f, 7, true) && drains == 1); // idempotent retirement + f.ngxWaitActive = true; f.ngxWaitCancelled = false; f.evaluatedFrames = 2; + assert(complete_ngx_input_wait(f, 7, false)); // output already drained: no extra sync + assert(drains == 1 && !f.ngxWaitActive && f.evaluatedFrames == 2); + f.ngxWaitActive = true; f.ngxWaitCancelled = false; + { std::lock_guard lock(f.ngxWaitMutex); + release_ngx_input_wait_locked(f, "timeout", true); } + assert(!f.ngxGpuWait); + drainResult = 700; + assert(!complete_ngx_input_wait(f, 7, true)); + assert(f.retiring && !f.ngxWaitActive && drains == 2); +} +''') + + def test_thread_clock_coarse_or_unavailable_reports_na(self): + bridge = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text() + clock = function(bridge, 'double WINAPI d4rThreadCpuTimeMs(void)') + run(r''' +#include +#include +#define WINAPI +long resolutionNs = 1; +bool available = true; +int fake_resolution(clockid_t, timespec* out) { out->tv_sec = 0; out->tv_nsec = resolutionNs; return 0; } +int timestamp(clockid_t, timespec* out) { out->tv_sec = 2; out->tv_nsec = 123456789; return available ? 0 : -1; } +#define clock_getres fake_resolution +#define clock_gettime timestamp +''' + clock + r''' +int main() { + assert(d4rThreadCpuTimeMs() == 2123.456789); + resolutionNs = 10000000; assert(d4rThreadCpuTimeMs() == -1); + resolutionNs = 1; available = false; assert(d4rThreadCpuTimeMs() == -1); +} +''') diff --git a/tests/test_latency_optimizations.py b/tests/test_latency_optimizations.py new file mode 100644 index 0000000..b6bfbe8 --- /dev/null +++ b/tests/test_latency_optimizations.py @@ -0,0 +1,204 @@ +"""CPU contracts for native completion and packed-output capability gating.""" +from pathlib import Path +import subprocess +import tempfile +import unittest + +ROOT = Path(__file__).resolve().parents[1] + + +def function(source, signature): + start = source.index(signature) + opening = source.index('{', start) + depth, end = 1, opening + 1 + while depth: + depth += (source[end] == '{') - (source[end] == '}') + end += 1 + return source[start:end] + + +def run(source, cpp=False): + with tempfile.TemporaryDirectory() as tmp: + exe = Path(tmp) / 'check' + result = subprocess.run(['c++' if cpp else 'cc', '-x', 'c++' if cpp else 'c', + '-std=c++20' if cpp else '-std=c11', '-pthread', '-Wall', '-Wextra', '-Werror', + '-I', str(ROOT / 'tools'), '-o', str(exe), '-'], + input=source, text=True, capture_output=True) + if result.returncode: raise AssertionError(result.stderr) + subprocess.run([str(exe)], check=True, timeout=5, cwd=tmp) + + +class LatencyContracts(unittest.TestCase): + def test_output_membership_filter_keeps_registry_results_exact(self): + bridge = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text() + start = bridge.index('enum { OUTPUT_KERNEL_CAPACITY') + end = bridge.index('/* 1 when the last output kernel', start) + code = bridge[start:end] + run(r''' +#include +#include +#include +#include +#include +#include +#include +#define CUfunction void* +#define R_OK 4 +static const char* fake_getenv(const char* name) { (void)name; return "/objects"; } +static int fake_access(const char* path, int mode) { (void)path; (void)mode; return 0; } +#define getenv fake_getenv +#define access fake_access +typedef struct { char name[128]; int rgba8; } NativeKernel; +NativeKernel objects[1] = {{"hiluma_engine_output_candidate", 1}}; +NativeKernel* native_kernels = objects; +size_t native_kernel_count = 1; +''' + code + r''' +int main(void) { + void* output = (void*)(uintptr_t)0x5000; + note_function_lookup(output, "hiluma_engine_output_candidate"); note_launch(output); + assert(last_output_native == 1 && last_output_rgba8 == 1); + for (uintptr_t i = 1; i < 1000; ++i) { + void* network = (void*)(i * 16); + if (network == output) continue; + note_launch(network); + assert(last_output_native == 1 && last_output_rgba8 == 1); // negatives/collisions change nothing + } + note_function_lookup(output, "rrlite_downsample_kernel_static_ldr"); note_launch(output); + assert(last_output_native == 1 && last_output_rgba8 == 0); // newest lookup wins on handle reuse +} +''') + + def test_completed_prefix_callback_gate(self): + run(r''' +#include "d4r_frame_completion.h" +#include +int main() { + D4rFrameCompletion completion; + assert(!completion.can_signal(1)); + completion.admit(1); completion.admit(2); completion.admit(3); + assert(completion.can_signal(1) && !completion.can_signal(2)); + assert(completion.retire(3) == 0); // later discard does not cross unfinished work + assert(!completion.can_signal(3)); + assert(completion.retire(1) == 1 && completion.can_signal(2)); + assert(completion.retire(2) == 3 && !completion.can_signal(3)); +} +''', cpp=True) + + def test_actual_native_callback_cancellation_monotonicity_and_detached_cleanup(self): + bridge = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text() + start = bridge.index('typedef struct', bridge.index("/* NGX's null-stream completion")) + end = bridge.index('} D4rNgxTimeline;', start) + len('} D4rNgxTimeline;') + definitions = bridge[start:end] + bodies = '\n'.join(function(bridge, sig) for sig in ( + 'int WINAPI d4rNgxTimelineSignal(', 'static void d4r_ngx_completed(', + 'CUresult WINAPI d4rNgxTimelinePrepare(', 'void WINAPI d4rNgxTimelineCancel(', + 'CUresult WINAPI d4rNgxTimelineQueue(', 'void WINAPI d4rNgxTimelineDestroy(', + 'CUresult WINAPI d4rNgxTimelineWait(')) + run(r''' +#define _POSIX_C_SOURCE 200809L +#include +#include +#include +#include +#include +#include +#define WINAPI +#define CUDA_SUCCESS 0 +#define CUDA_ERROR_NOT_READY 600 +#define CUDA_ERROR_NOT_SUPPORTED 801 +#define CUDA_ERROR_UNKNOWN 999 +#define CUresult int +typedef struct { int sType; const void* pNext; uint64_t semaphore, value; } D4rVkSemaphoreSignalInfo; +typedef int (*VK_SIGNAL_SEMAPHORE_FN)(void*, const D4rVkSemaphoreSignalInfo*); +typedef int (*HIP_LAUNCH_HOST_FN)(void*, void(*)(void*), void*); +static void* d4r_ngx_stream(void* stream) { return stream; } /* D4R_NGX_STREAM_PRIORITY off */ +void (*callback)(void*); void* argument; +uint64_t signalled = 0; int signals = 0, launchResult = 0; +int signal_fake(void* device, const D4rVkSemaphoreSignalInfo* info) { + (void)device; assert(info->value > signalled); signalled = info->value; ++signals; return 0; +} +static void* complete_later(void* unused) { + (void)unused; struct timespec pause = {0, 20000000}; nanosleep(&pause, NULL); + callback(argument); return NULL; +} +int launch_fake(void* stream, void(*fn)(void*), void* arg) { + assert(stream == NULL); + if (!launchResult) { callback = fn; argument = arg; } + return launchResult; +} +''' + definitions + '\n' + bodies + r''' +int main(void) { + D4rNgxTimeline* t = calloc(1, sizeof(*t)); + assert(pthread_mutex_init(&t->lock, NULL) == 0); + pthread_condattr_t attr; assert(pthread_condattr_init(&attr) == 0); + assert(pthread_condattr_setclock(&attr, CLOCK_MONOTONIC) == 0); + assert(pthread_cond_init(&t->completed, &attr) == 0); pthread_condattr_destroy(&attr); + t->signal = signal_fake; t->launch = launch_fake; + assert(d4rNgxTimelinePrepare(t) == 0); + assert(d4rNgxTimelineQueue(t, 1) == 0); + assert(d4rNgxTimelinePrepare(t) == 600 && d4rNgxTimelineQueue(t, 2) == 600); + d4rNgxTimelineCancel(t); callback(argument); assert(signals == 0 && t->busy == 0); + assert(d4rNgxTimelineWait(t) == 0); + assert(d4rNgxTimelineSignal(t, 1) == 0 && signals == 1); + assert(d4rNgxTimelinePrepare(t) == 0 && d4rNgxTimelineQueue(t, 2) == 0); + callback(argument); assert(signalled == 2 && signals == 2); + assert(d4rNgxTimelinePrepare(t) == 0 && d4rNgxTimelineQueue(t, 3) == 0); + assert(d4rNgxTimelineSignal(t, 4) == 0); callback(argument); + assert(signalled == 4 && signals == 3); // late callback never signals backwards + launchResult = 1; + assert(d4rNgxTimelinePrepare(t) == 0 && d4rNgxTimelineQueue(t, 5) == 801 && t->busy == 0); + launchResult = 0; + assert(d4rNgxTimelineQueue(t, 5) == 0); + pthread_t worker; assert(pthread_create(&worker, NULL, complete_later, NULL) == 0); + assert(d4rNgxTimelineWait(t) == 0); pthread_join(worker, NULL); + assert(signalled == 5 && signals == 4); + assert(d4rNgxTimelinePrepare(t) == 0 && d4rNgxTimelineQueue(t, 6) == 0); + d4rNgxTimelineDestroy(t); callback(argument); // ticket outlives owner, cancellation forbids a signal + assert(signals == 4); +} +''') + + def test_all_half_values_pack_like_the_reference_unorm_conversion(self): + header = (ROOT / 'kernels/tex/tex_common.h').read_text() + body = function(header, '__attribute__((device)) static inline uint32_t redirect_unorm(') + body = body.replace('__attribute__((device)) ', '') + run(r''' +#include +#include +#include +''' + body + r''' +int main() { + for (uint32_t bits = 0; bits < 65536; ++bits) { + const int exponent = (bits >> 10) & 31, mantissa = bits & 1023; + double value = exponent == 0 ? std::ldexp(double(mantissa), -24) + : exponent == 31 ? (mantissa ? NAN : INFINITY) + : std::ldexp(1.0 + double(mantissa) / 1024.0, exponent - 15); + if (bits & 0x8000) value = -value; + value = value > 0.0 ? value : 0.0; + value = value < 1.0 ? value : 1.0; + const uint32_t expected = uint32_t(std::floor(value * 255.0 + 0.5)); + assert(redirect_unorm(uint16_t(bits), 255.0f) == expected); + } +} +''', cpp=True) + + def test_packed_capability_requires_matching_code_object_bytes(self): + bridge = (ROOT / 'tools/wine_nvcuda_bridge.c').read_text() + body = function(bridge, 'static int native_rgba8_verified(') + run(r''' +#include +#include +#include +#include +''' + body + r''' +int main(void) { + assert(!native_rgba8_verified("/does/not/exist", "output")); + FILE* f = fopen("output.hsaco", "wb"); assert(f != NULL); fputs("abc", f); fclose(f); + assert(!native_rgba8_verified(".", "output")); + f = fopen("d4r-rgba8.txt", "w"); fputs("output e71fa2190541574b\n", f); fclose(f); + assert(native_rgba8_verified(".", "output")); + f = fopen("output.hsaco", "ab"); fputc('d', f); fclose(f); + assert(!native_rgba8_verified(".", "output")); + assert(!native_rgba8_verified(".", "other")); +} +''') diff --git a/tests/test_shim_config.py b/tests/test_shim_config.py index a7b9e6e..e174bba 100644 --- a/tests/test_shim_config.py +++ b/tests/test_shim_config.py @@ -35,10 +35,13 @@ class ShimConfigTests(unittest.TestCase): ini = Path(temporary) / 'd4r.ini' env = {key: value for key, value in os.environ.items() if not key.startswith('D4R_')} cmd = ['python3', str(ROOT / 'scripts/d4r_config.py'), '--config', str(ini)] - for value, expected in ((None, '200'), ('auto', '200'), ('200', '200'), ('0', '0'), ('75', '75')): + for value, expected in ((None, None), ('auto', None), ('200', '200'), ('0', '0'), ('75', '75')): ini.write_text('[Interop]\n' + (f'MarkerPollUs = {value}\n' if value is not None else '')) result = subprocess.run(cmd, env=env, text=True, capture_output=True, check=True) - self.assertIn(f'export D4R_SHIM_MARKER_POLL_US={expected}\n', result.stdout) + if expected is None: + self.assertNotIn('export D4R_SHIM_MARKER_POLL_US=', result.stdout) + else: + self.assertIn(f'export D4R_SHIM_MARKER_POLL_US={expected}\n', result.stdout) # An explicit environment variable keeps its existing precedence. result = subprocess.run(cmd, env=dict(env, D4R_SHIM_MARKER_POLL_US='35'), text=True, capture_output=True, check=True) diff --git a/tools/d3d12_dlss_harness.cpp b/tools/d3d12_dlss_harness.cpp index 982e91b..489eaa6 100644 --- a/tools/d3d12_dlss_harness.cpp +++ b/tools/d3d12_dlss_harness.cpp @@ -1125,9 +1125,9 @@ int main(int argc, char** argv) std::fprintf(stderr, "RGB10A2 output cannot be combined with RGBA8 or the quality scene\n"); return 2; } - if (rgba8 && (qualityScene || motionScene || jitterScene || replayDir != nullptr)) + if (rgba8 && (qualityScene || replayDir != nullptr)) { - std::fprintf(stderr, "RGBA8 format probe requires the static synthetic scene\n"); + std::fprintf(stderr, "RGBA8 format probe supports static, motion and jitter scenes\n"); return 2; } const bool temporal = std::getenv("D4R_HARNESS_TEMPORAL") != nullptr || motionScene || jitterScene || qualityScene || @@ -1226,6 +1226,8 @@ int main(int argc, char** argv) for (size_t index = 0; index < color.size(); ++index) color8[index] = static_cast(std::lround( std::clamp(static_cast(std::bit_cast<_Float16>(color[index])), 0.0f, 1.0f) * 255.0f)); + if (std::getenv("D4R_HARNESS_RGBA8_RAMP") != nullptr) + for (size_t i = 0; i < color8.size(); ++i) color8[i] = static_cast(i); outputInit8.resize(static_cast(outWidth) * outHeight * 4); } @@ -1421,6 +1423,17 @@ int main(int argc, char** argv) : rgb10 ? DXGI_FORMAT_R10G10B10A2_UNORM : DXGI_FORMAT_R16G16B16A16_FLOAT, D3D12_RESOURCE_FLAG_ALLOW_UNORDERED_ACCESS); + // Moving/jittered RGBA8 fixtures quantize the same canonical half input as the initial upload. + // Never upload the half byte layout into an RGBA8 resource: that would make a format test silently + // exercise different content instead of checking the conversion/packing route used by the game. + const auto update_color = [&] { + if (rgba8) + for (size_t i = 0; i < color.size(); ++i) + color8[i] = static_cast(std::lround(std::clamp( + static_cast(std::bit_cast<_Float16>(color[i])), 0.0f, 1.0f) * 255.0f)); + update_texture(colorTexture, rgba8 ? static_cast(color8.data()) + : static_cast(color.data()), resourceWidth * (rgba8 ? 4 : 8), srv); + }; upload(colorTexture, rgba8 ? static_cast(color8.data()) : static_cast(color.data()), resourceWidth * (rgba8 ? 4 : 8), srv); upload(depthTexture, depth.data(), resourceWidth * 4, srv); @@ -1757,7 +1770,7 @@ int main(int argc, char** argv) color[pixel * 4 + 0] = color[pixel * 4 + 1] = color[pixel * 4 + 2] = half; color[pixel * 4 + 3] = 0x3c00u; } - update_texture(colorTexture, color.data(), resourceWidth * 8, srv); + update_color(); if (frame == 1 && (panX != 0 || panY != 0)) { // Current position -> previous position. Without MVLowRes DLSS @@ -1800,7 +1813,7 @@ int main(int argc, char** argv) {0.0f, 0.25f, -0.25f, -0.125f, 0.375f, 0.125f, -0.375f, 0.0f}; const size_t phase = static_cast(frame - 1) % jitterX.size(); render_jittered_background(color, inWidth, inHeight, resourceWidth, jitterX[phase], jitterY[phase]); - update_texture(colorTexture, color.data(), resourceWidth * 8, srv); + update_color(); d4r_ngx_set_float(parameters, "Jitter.Offset.X", jitterX[phase]); d4r_ngx_set_float(parameters, "Jitter.Offset.Y", jitterY[phase]); } @@ -1823,7 +1836,7 @@ int main(int argc, char** argv) motion[motionPixel * 2 + 0] = 0xc400u; // -4 render pixels: current position -> previous motion[motionPixel * 2 + 1] = 0; } - update_texture(colorTexture, color.data(), resourceWidth * 8, srv); + update_color(); update_texture(depthTexture, depth.data(), resourceWidth * 4, srv); update_texture(motionTexture, motion_rows(motion.data()), motionRowBytes, srv); } @@ -2054,6 +2067,20 @@ int main(int argc, char** argv) LARGE_INTEGER counterStart, counterRecorded, counterEnd, counterFrequency; QueryPerformanceFrequency(&counterFrequency); QueryPerformanceCounter(&counterStart); + // D4R_HARNESS_THRASH_MB=N: copy N MB between two private buffers ahead of the evaluation, as a game's + // own rendering does, so DLSS starts with its weights and history evicted from L2 and Infinity Cache + // instead of finding them warm from the previous back-to-back frame. + static const UINT64 thrashBytes = std::getenv("D4R_HARNESS_THRASH_MB") != nullptr + ? std::clamp(std::strtoull(std::getenv("D4R_HARNESS_THRASH_MB"), nullptr, 10), 0, 1024) << 20 : 0; + static ID3D12Resource* thrashSource = nullptr; + static ID3D12Resource* thrashTarget = nullptr; + if (thrashBytes != 0 && thrashSource == nullptr) + { + thrashSource = create_buffer(D3D12_HEAP_TYPE_DEFAULT, thrashBytes, D3D12_RESOURCE_STATE_COMMON); + thrashTarget = create_buffer(D3D12_HEAP_TYPE_DEFAULT, thrashBytes, D3D12_RESOURCE_STATE_COMMON); + } + if (thrashSource != nullptr && thrashTarget != nullptr) + g_list->CopyBufferRegion(thrashTarget, 0, thrashSource, 0, thrashBytes); result = evaluate(g_list, feature, parameters, nullptr); QueryPerformanceCounter(&counterRecorded); if (const char* delay = std::getenv("D4R_HARNESS_DELAY_SUBMIT_MS")) diff --git a/tools/d4r_frame_completion.h b/tools/d4r_frame_completion.h index 12a7f94..385d561 100644 --- a/tools/d4r_frame_completion.h +++ b/tools/d4r_frame_completion.h @@ -11,6 +11,11 @@ class D4rFrameCompletion public: void admit(uint32_t frame) { pending_.push_back({frame, false}); } + // A native completion callback can release this value without crossing an unfinished producer + // only when it is the first admitted frame still outstanding. Retirement itself stays on the + // worker, which also folds any later discarded frames into the completed prefix. + bool can_signal(uint32_t frame) const { return !pending_.empty() && pending_.front().frame == frame; } + uint64_t retire(uint32_t frame) { for (auto& pending : pending_) diff --git a/tools/d4r_input_wait.h b/tools/d4r_input_wait.h new file mode 100644 index 0000000..5f13312 --- /dev/null +++ b/tools/d4r_input_wait.h @@ -0,0 +1,30 @@ +#pragma once +#include + +// The original marker remains the authority for whether a command list ran. A CPU release of the +// separate, coherent GPU wait word must never make discarded input look ready to the worker. +enum class D4rInputWaitAbort { None, Retired, Discarded, Timeout }; + +inline D4rInputWaitAbort d4r_input_wait_abort(uint32_t frame, uint32_t slotMarker, uint32_t aggregate, + bool retiring, double elapsedUs) +{ + if (retiring) + return D4rInputWaitAbort::Retired; + if (slotMarker != frame && static_cast(aggregate - frame) > 0) + return D4rInputWaitAbort::Discarded; + // Also covers a stream wait that failed to observe an input word already written by Vulkan. + if (elapsedUs > 5000000.0) + return D4rInputWaitAbort::Timeout; + return D4rInputWaitAbort::None; +} + +inline unsigned d4r_marker_poll_delay_us(double predictedUs, double elapsedUs) +{ + const double remaining = predictedUs - 150.0 - elapsedUs; + if (remaining > 20.0) + return static_cast(remaining < 200.0 ? remaining : 200.0); + if (elapsedUs < predictedUs + 150.0) + return 0; // yield during the final 150 us and a short late window + const double late = 20.0 * (1.0 + (elapsedUs - predictedUs - 150.0) / 160.0); + return static_cast(late < 200.0 ? late : 200.0); +} diff --git a/tools/d4r_nvngx_shim.cpp b/tools/d4r_nvngx_shim.cpp index debc5fa..f94b01a 100644 --- a/tools/d4r_nvngx_shim.cpp +++ b/tools/d4r_nvngx_shim.cpp @@ -30,6 +30,15 @@ // D4R_SHIM_INPUT_STATE / D4R_SHIM_DEPTH_STATE / D4R_SHIM_OUTPUT_STATE // D3D12 resource states the inputs/output are in when // EvaluateFeature is called (defaults 0x40, 0x40, 0x8) +// D4R_SHIM_GPU_INPUT_WAIT=1 experimental NGX null-stream input wait (default 0; split VRAM only) +// D4R_SHIM_OUTPUT_RGBA8=1 validated native K output stores RGBA8 directly (default 0) +// D4R_SHIM_NATIVE_COMPLETION=1 native null-stream callback signals/wakes split completion (default 0) +// D4R_SHIM_FUSE_RGBA8_INPUT=1 combine RGBA8 color conversion and linear-buffer copy (default 0) +// D4R_SHIM_GPU_COMPLETION=1 (needs GPU_INPUT_WAIT) the game's queue polls DLSS's completion word on the GPU +// instead of a CPU-signalled split (bounded poll; default 0) +// D4R_PROFILE_GPU_EVENTS / _TIMESTAMPS / _LOG / _CPU_TIME=0 drop one part of D4R_PROFILE (bisection) +// D4R_NGX_STREAM_PRIORITY=high|normal (bridge) run NGX on one priority HIP stream (default off; no gain measured) +// D4R_SHIM_MARKER_POLL_US=N fixed CPU marker pacing; unset: adaptive (0 yields) // D4R_SHIM_VRAM_INTEROP=1 keep inputs and output in VRAM (see "VRAM interop") // D4R_SHIM_SPLIT_FRAME=1 with VRAM interop and the d4r vkd3d-proton patch, // present each frame's own DLSS result (see "Split frames") @@ -54,6 +63,8 @@ #include #include #include "d4r_motion_dilation.h" +#include "d4r_input_wait.h" +#include "d4r_vulkan_staging.h" #include #include @@ -266,10 +277,23 @@ struct CudaApi int(WINAPI* memcpy2DAsync)(const CudaMemcpy2D*, void*) = nullptr; // Optional (d4r nvcuda bridge): null-stream wait for a device u32, and a write from another stream. int(WINAPI* streamWaitValue32)(CudaDevicePtr, uint32_t) = nullptr; + int(WINAPI* streamWaitValue32Eq)(CudaDevicePtr, uint32_t) = nullptr; + int(WINAPI* ngxWriteValue32)(CudaDevicePtr, uint32_t) = nullptr; // D4R_SHIM_GPU_COMPLETION + int(WINAPI* ngxArmCompletion)(CudaDevicePtr, uint32_t, int) = nullptr; // early word behind the output kernel + int(WINAPI* ngxCompletionFired)() = nullptr; + double(WINAPI* threadCpuTimeMs)() = nullptr; int(WINAPI* writeValue32)(CudaDevicePtr, uint32_t) = nullptr; // Optional (d4r nvcuda bridge): surfaces of an array store to linear memory instead (d4r native kernels). int(WINAPI* setArrayRedirect)(CudaArray, CudaDevicePtr, uint32_t) = nullptr; int(WINAPI* outputKernelNative)() = nullptr; + int(WINAPI* outputKernelRgba8)() = nullptr; + int(WINAPI* timelineCreate)(VkDevice, uint64_t, void**) = nullptr; + int(WINAPI* timelinePrepare)(void*) = nullptr; + void(WINAPI* timelineCancel)(void*) = nullptr; + int(WINAPI* timelineQueue)(void*, uint64_t) = nullptr; + int(WINAPI* timelineSignal)(void*, uint64_t) = nullptr; + void(WINAPI* timelineDestroy)(void*) = nullptr; + int(WINAPI* timelineWait)(void*) = nullptr; // Optional (d4r nvcuda bridge): report a pitch-linear texture to NGX as an array of that size/format. int(WINAPI* registerLinearTexture)(CudaObject, size_t, size_t, uint32_t, uint32_t) = nullptr; int(WINAPI* dilateMotion)(const D4rMotionDilationParams*) = nullptr; @@ -665,7 +689,8 @@ static void load_portable_config() portable_set(setting.variable, ini_flag(ini, "interop", setting.key, setting.fallback) ? "1" : "0"); portable_set_unix("D4R_ELIDE_NGX_SYNC", ini_flag(ini, "interop", "ElideNgxSync", 1) ? "1" : "0"); const std::string poll = ini_value(ini, "interop", "MarkerPollUs"); - portable_set("D4R_SHIM_MARKER_POLL_US", poll.empty() ? "200" : poll); + if (!poll.empty() && ascii_lower(poll) != "auto") + portable_set("D4R_SHIM_MARKER_POLL_US", poll); // Direct output takes effect only while NGX runs the native output kernel (the bridge reports it); the // release cannot ship that kernel, which is built from NVIDIA's PTX. portable_set("D4R_SHIM_OUTPUT_DIRECT", ini_flag(ini, "interop", "DirectOutput", nativeOn) ? "1" : "0"); @@ -786,16 +811,6 @@ static double profile_ms(ProfileClock::time_point start, ProfileClock::time_poin return std::chrono::duration(end - start).count(); } -static double profile_thread_cpu_ms() -{ - FILETIME created = {}, exited = {}, kernel = {}, user = {}; - if (!GetThreadTimes(GetCurrentThread(), &created, &exited, &kernel, &user)) - return -1.0; - const uint64_t kernelTicks = (static_cast(kernel.dwHighDateTime) << 32) | kernel.dwLowDateTime; - const uint64_t userTicks = (static_cast(user.dwHighDateTime) << 32) | user.dwLowDateTime; - return static_cast(kernelTicks + userTicks) / 10000.0; -} - // --- CUDA worker thread -------------------------------------------------------- // Every NGX-core and CUDA call runs on this thread so the ZLUDA context that // the nvcuda bridge creates on first use stays current. @@ -1357,6 +1372,14 @@ static bool denoiser_has_preset_f() { return g.denoiserVersion[0] > 310 || (g.denoiserVersion[0] == 310 && g.denoiserVersion[1] >= 9); } + +static double profile_thread_cpu_ms() +{ + // Wine's GetThreadTimes is tick-accounted (10 ms on this host), despite FILETIME's 100 ns units. + // The Unix bridge reads the current worker's actual thread CPU clock. Older bridges report n/a. + static const bool wanted = env_uint("D4R_PROFILE_CPU_TIME", 1) != 0; // 0: skip the Unix call (bisection) + return wanted && g.cu.threadCpuTimeMs != nullptr ? g.cu.threadCpuTimeMs() : -1.0; +} static std::mutex g_featureApiMutex; template static bool load_export(HMODULE module, const char* name, Function& function) @@ -1536,10 +1559,31 @@ static bool load_libraries() if (!load_export(g.cuda, "d4rStreamWaitValue32", g.cu.streamWaitValue32) || !load_export(g.cuda, "d4rWriteValue32", g.cu.writeValue32)) g.cu.streamWaitValue32 = nullptr, g.cu.writeValue32 = nullptr; + g.cu.streamWaitValue32Eq = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rStreamWaitValue32Eq"))); + g.cu.ngxWriteValue32 = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxWriteValue32"))); + g.cu.ngxArmCompletion = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxArmCompletion"))); + g.cu.ngxCompletionFired = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxCompletionFired"))); + if (g.cu.ngxArmCompletion == nullptr || g.cu.ngxCompletionFired == nullptr) + g.cu.ngxArmCompletion = nullptr, g.cu.ngxCompletionFired = nullptr; + g.cu.threadCpuTimeMs = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rThreadCpuTimeMs"))); if (!load_export(g.cuda, "d4rSetArrayRedirect", g.cu.setArrayRedirect)) g.cu.setArrayRedirect = nullptr; if (!load_export(g.cuda, "d4rOutputKernelNative", g.cu.outputKernelNative)) g.cu.outputKernelNative = nullptr; + g.cu.outputKernelRgba8 = reinterpret_cast( + reinterpret_cast(GetProcAddress(g.cuda, "d4rOutputKernelRgba8"))); + g.cu.timelineCreate = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineCreate"))); + g.cu.timelinePrepare = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelinePrepare"))); + g.cu.timelineCancel = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineCancel"))); + g.cu.timelineQueue = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineQueue"))); + g.cu.timelineSignal = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineSignal"))); + g.cu.timelineDestroy = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineDestroy"))); + g.cu.timelineWait = reinterpret_cast(reinterpret_cast(GetProcAddress(g.cuda, "d4rNgxTimelineWait"))); if (!load_export(g.cuda, "d4rRegisterLinearTexture", g.cu.registerLinearTexture)) g.cu.registerLinearTexture = nullptr; g.cu.dilateMotion = reinterpret_cast( @@ -1878,6 +1922,15 @@ struct VramBuffer uint8_t kind = 0; }; +// Per-slot descriptors and their source texture live until recorded allocators retire. +struct FusedColorInput { + VkImage image = VK_NULL_HANDLE; + VkImageView view = VK_NULL_HANDLE; + VkBuffer buffer = VK_NULL_HANDLE; + D4rVkStagingSet set; + ID3D12Resource* owner = nullptr; +}; + // A temporary image for colour or exposure format conversion. struct VramImage { @@ -2318,6 +2371,8 @@ struct FrameParams int createQuality = 0, createFlags = 0, createSubrects = 0; unsigned int createPreset = 0; bool vram = false; // inputs and output stay in VRAM (VRAM interop) + bool gpuInputMarker = false; // this list recorded the exportable per-slot readiness mirror + bool gpuCompletion = false; // the game's queue polls this frame's completion word instead of a split wait bool split = false; // this frame presents its own result (split frames) unsigned motionDilation = 0; // actually applied, worker only // D4R_SHIM_INPLACE: the game's linear output texture, written without a game-side copy (0 = none) @@ -2326,6 +2381,7 @@ struct FrameParams // D4R_SHIM_OUTPUT_PACKED: the output kernel stores the game's R10G10B10A2 texels itself, so the game-side // copy (or the in-place texture) takes them without the RGBA16F conversion image and blit bool packedOutput = false; + bool packedRgba8 = false; // R8G8B8A8_UNORM after the original half rounding // Ray Reconstruction: the denoiser's own per-frame settings and camera matrices, copied // off the caller's parameter object on the game thread so the worker never reads an // object the game may have destroyed by then. A scalar the caller did not set this frame @@ -2350,11 +2406,23 @@ struct FrameTiming double download = 0, outputConvert = 0; double prep = 0, worker = 0, workerWait = 0, finishWait = 0, markerWait = 0; double ngxHost = 0, ctxSync = 0, gpuEval = -1; + // GPU-input-wait frames split gpu_eval at an event recorded right behind the wait: input_gpu_wait is launch -> + // inputs released (the game's remaining input work), dlss_gpu is inputs released -> completion (the real DLSS + // cost). CPU-wait frames have no wait on the GPU: dlss_gpu equals gpu_eval and input_gpu_wait is 0. + double dlssGpu = -1, inputGpuWait = -1; + // Game-queue GPU timestamps around this evaluation (D3D12 queries, read 4 frames later, so these describe the + // frame kOutputSlots earlier): vk_input = shim input commands, vk_dlss_wait = input marker -> second half + // starts (input release, DLSS, completion signal), vk_output = second-half commands, vk_total = all of it + // (what an overlay timing the evaluate call on the game's list reports). + double vkInput = -1, vkDlssWait = -1, vkOutput = -1, vkTotal = -1; + bool inputEventRecorded = false; double motionDilationWall = 0; bool markerDeferred = false; // the worker waits for the input marker (D4R_SHIM_DEFER_MARKER) + bool gpuInputWait = false; // GPU events span wait enqueue -> NGX completion, including input readiness bool gpuEventsRecorded = false; double outputSyncWall = 0, outputSyncCpu = -1; bool outputSyncBlocking = false; + bool outputSyncNative = false; double h2dTotal = 0, d2hIssue = 0, d2hWait = 0; ProfileClock::time_point prepStart, workerQueued, finishQueued; size_t queueDepth = 0; @@ -2512,12 +2580,20 @@ struct Feature // Output.Subrect origins are read at all, so a replay must create the feature the same way. int outputSubrects = 0; CudaEvent profileStart = nullptr, profileEnd = nullptr; + CudaEvent profileInputsReleased = nullptr; // D4R_PROFILE only: recorded right behind the GPU-side input wait + // D4R_PROFILE only: 4 game-queue timestamps per output slot (start, input marker, second half, end) + ID3D12QueryHeap* timestampHeap = nullptr; + ID3D12Resource* timestampReadback = nullptr; + volatile uint64_t* timestampValues = nullptr; + double timestampTicksPerMs = 0.0; + bool timestampFailed = false; CudaEvent outputReadyEvent = nullptr; // output wait pacing (D4R_SHIM_OUTPUT_POLL_US): smoothed time from the event record to its completion, kept // separately for each wait site (they wait for very different amounts of GPU work) double publishWaitEmaUs = 0.0, evalWaitEmaUs = 0.0; // native Vulkan input wait: smoothed time from recording to the input event (waiters may overlap) std::atomic vkInputWaitEmaUs{0.0}; + std::atomic markerWaitEmaUs{0.0}; // prep or worker waits; never both for the same frame bool profileEventsReady = false; CudaImage color, depth, motion, exposure, output; @@ -2541,6 +2617,9 @@ struct Feature // VRAM interop: decided on the first evaluation, dropped for good if the // formats later stop qualifying. bool vram = false, vramDecided = false; + bool fusedColorAllowed = true; + FusedColorInput fusedColor[kSlots]; + std::vector retiredFusedColor; VramImage colorConversion; VramImage depthConversion; VramImage motionConversion; @@ -2553,14 +2632,14 @@ struct Feature std::mutex splitMutex; uint64_t splitSignalled = 0; // under splitMutex D4rFrameCompletion splitCompletion; // under splitMutex; cancellation can finish out of order - // Inputs become ready on the CPU before CUDA work is queued. GPU-side waits - // for recorded-but-discarded command lists can otherwise wedge the device. + // GPU input waits are opt-in, with an independent CPU watchdog that can release discarded lists. // D4R_SHIM_OUTPUT_DIRECT: this frame's result is stored by the native output kernel straight into its // destination buffer (worker only), so the array -> buffer copy is skipped. bool outputRedirected = false; // The previous frame's output kernel was the native one and honoured the redirect (worker writes, game // thread reads): only then is a frame recorded for packed output. std::atomic packedReady{false}; + std::atomic rgba8Ready{false}; bool outputDirectAllowed = false; // the forced preset's output kernel honours the redirect // D4R_SHIM_LINEAR_INPUTS: NGX samples the interop buffers directly through pitch-linear // texture objects (rows padded to 256 bytes), so the buffer -> array copies disappear. @@ -2570,6 +2649,24 @@ struct Feature std::vector retiredImages; ID3D12Resource* marker = nullptr; volatile uint32_t* markerValue = nullptr; // [0] latest input; [1 + slot] exact input frame + // The D3D12 READBACK marker has no exportable HIP pointer. Mirror its per-slot tags in a dedicated + // HOST_VISIBLE | HOST_COHERENT Vulkan allocation imported into HIP, as engineSync is. Keep the original + // marker as the authority: releasing a wait from the CPU must not turn unsubmitted input into ready input. + VramBuffer ngxInputSync; + // D4R_SHIM_GPU_COMPLETION: words kCompletionWord + frame % kOutputSlots of ngxInputSync carry each frame's + // completion (written by HIP after the output, or by the CPU in frame_retired); the game's queue polls them. + bool gpuCompletion = false; + bool completionFired = false; // worker only: the bridge queued this frame's word behind the native output kernel + D4rVkStagingSet completionSet; + void* ngxTimeline = nullptr; // native completion and CPU drops share this timeline writer + bool ngxNativeCompletion = false; // worker only after feature setup + std::atomic ngxGpuWait{false}; + bool ngxWaitEventsWarmed = false; // worker only: initialize event storage before the first stream wait + std::mutex ngxWaitMutex; + bool ngxWaitActive = false, ngxWaitCancelled = false; // under ngxWaitMutex, armed before calling HIP + uint32_t ngxWaitFrame = 0; + int ngxWaitSlot = 0; + ProfileClock::time_point ngxWaitStart; uint32_t frame = 0; uint32_t rejectedFormatCount = 0; // game thread; rate-limits format diagnostics std::atomic latestOutput{-1}; @@ -2672,6 +2769,14 @@ static bool create_buffer(D3D12_HEAP_TYPE heapType, UINT64 size, ID3D12Resource* return true; } +// D4R_PROFILE: game-queue timestamps for one evaluation. Point 0 is written before the input commands, 1 after the +// input marker, 2 at the start of the split's second half, 3 after the output commands (which also resolves the +// slot). Reading slot frame % kOutputSlots returns the evaluation kOutputSlots frames earlier, whose second half has +// normally run by now; incomplete (zero) values are skipped. Allocation failure disables it for the feature. +struct Feature; +static bool profile_timestamps_ready(Feature& feature); +static void profile_timestamp(Feature& feature, ID3D12GraphicsCommandList* list, uint32_t frame, int point); + // Ensures staging matches the resource's first subresource; heap selects // readback (inputs) or upload (output). static bool ensure_staging(Staging& staging, ID3D12Resource* resource, D3D12_HEAP_TYPE heapType) @@ -3036,6 +3141,8 @@ struct VulkanInterop PFN_vkDestroyBuffer destroyBuffer = nullptr; PFN_vkGetBufferMemoryRequirements bufferRequirements = nullptr; PFN_vkBindBufferMemory bindBuffer = nullptr; + PFN_vkCreateImageView createView = nullptr; + PFN_vkDestroyImageView destroyView = nullptr; PFN_vkCreateImage createImage = nullptr; PFN_vkDestroyImage destroyImage = nullptr; PFN_vkGetImageMemoryRequirements imageRequirements = nullptr; @@ -3056,6 +3163,14 @@ struct VulkanInterop int(WINAPI* release)(void*) = nullptr; }; static VulkanInterop g_vk; +static D4rVkStaging g_fusedColorStage; +// D4R_SHIM_GPU_COMPLETION: the bounded completion-word wait (kernels/support/vulkan_wait.comp) +static D4rVkStaging g_waitStage; +static bool g_waitReady = false; +constexpr uint32_t kCompletionWord = 16; +// ~1 us per poll of coherent memory: at most about a second, far below the kernel driver's 10 s GPU timeout +constexpr uint32_t kCompletionPollLimit = 1u << 20; +static bool g_fusedColorReady = false; static bool vram_interop_requested() { @@ -3117,6 +3232,8 @@ static void init_vram_interop() load(g_vk.destroyBuffer, "vkDestroyBuffer"); load(g_vk.bufferRequirements, "vkGetBufferMemoryRequirements"); load(g_vk.bindBuffer, "vkBindBufferMemory"); + load(g_vk.createView, "vkCreateImageView"); + load(g_vk.destroyView, "vkDestroyImageView"); load(g_vk.createImage, "vkCreateImage"); load(g_vk.destroyImage, "vkDestroyImage"); load(g_vk.imageRequirements, "vkGetImageMemoryRequirements"); @@ -3141,6 +3258,16 @@ static void init_vram_interop() FAILED(g.device->QueryInterface(__uuidof(ID3D12DXVKInteropDeviceD4R3), reinterpret_cast(&g_vk.linear)))) g_vk.linear = nullptr; g_vk.ready = ok; + if (ok && env_uint("D4R_SHIM_FUSE_RGBA8_INPUT", 0) != 0) + { + g_fusedColorReady = g_fusedColorStage.init(g_vk.device, deviceProc); + if (!g_fusedColorReady) g_fusedColorStage.destroy(); + } + if (ok && env_uint("D4R_SHIM_GPU_COMPLETION", 0) != 0) + { + g_waitReady = g_waitStage.init(g_vk.device, deviceProc) && g_waitStage.wait_ready(); + if (!g_waitReady) g_waitStage.destroy(); + } logf("VRAM interop: %s (VkDevice %p), split frames %s", ok ? "ready" : "missing Vulkan entry points", static_cast(g_vk.device), g_vk.split != nullptr ? "available" : "unavailable (stock vkd3d-proton)"); } @@ -3200,6 +3327,7 @@ static bool create_vram_buffer(VramBuffer& target, size_t bytes, bool importNow VkBufferCreateInfo info = {VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO, &external}; info.size = bytes; info.usage = VK_BUFFER_USAGE_TRANSFER_SRC_BIT | VK_BUFFER_USAGE_TRANSFER_DST_BIT; + if (g_fusedColorReady) info.usage |= VK_BUFFER_USAGE_STORAGE_BUFFER_BIT; info.sharingMode = VK_SHARING_MODE_EXCLUSIVE; VkResult result = g_vk.createBuffer(g_vk.device, &info, nullptr, &buffer.buffer); VkMemoryRequirements requirements = {}; @@ -3546,10 +3674,41 @@ static VkBufferImageCopy full_region(const VramCopy& copy) return region; } +static bool ensure_fused_color(Feature& feature, int slotIndex, ID3D12Resource* source, const VramCopy& copy) +{ + if (!g_fusedColorReady || !feature.fusedColorAllowed || env_uint("D4R_SHIM_FUSE_RGBA8_INPUT", 0) == 0 || source->GetDesc().Format != DXGI_FORMAT_R8G8B8A8_UNORM || + (source->GetDesc().Flags & D3D12_RESOURCE_FLAG_DENY_SHADER_RESOURCE) != 0) return false; + // Descriptor/image lifetimes extend through allocator retirement, including discarded lists. + // Bound that retention if a game creates a new transient input resource every frame; fallback + // to the original blit/copy path without invalidating descriptors already recorded on the GPU. + if (feature.retiredFusedColor.size() >= 64) + { + feature.fusedColorAllowed = false; + logf("fused color input descriptor retention reached its bound; using blit/copy for this feature"); + return false; + } + auto& slot = feature.fusedColor[slotIndex]; + auto& target = feature.inputs[slotIndex].vram[0]; + if (slot.image == copy.image && slot.buffer == target.buffer && slot.owner == source) return true; + if (slot.owner != nullptr) feature.retiredFusedColor.push_back(slot); + slot = {}; + VkImageViewCreateInfo view = {VK_STRUCTURE_TYPE_IMAGE_VIEW_CREATE_INFO}; + view.image = copy.image; view.viewType = VK_IMAGE_VIEW_TYPE_2D; view.format = VK_FORMAT_R8G8B8A8_UNORM; + view.subresourceRange = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1}; + if (g_vk.createView(g_vk.device, &view, nullptr, &slot.view) != VK_SUCCESS) return false; + if (!g_fusedColorStage.set(slot.set, slot.view, view.format, target.buffer, feature.inputs[slotIndex].host[0].size(), false)) + { + g_fusedColorStage.free(slot.set); g_vk.destroyView(g_vk.device, slot.view, nullptr); slot = {}; return false; + } + slot.image = copy.image; slot.buffer = target.buffer; slot.owner = source; source->AddRef(); + logf("fused RGBA8 input conversion active for slot %d", slotIndex); + return true; +} + // Records the input copies into the slot's buffers. The resources are in // COPY_SOURCE state already. static bool record_vram_inputs(Feature& feature, ID3D12GraphicsCommandList* list, InputSlot& slot, - const VramCopy* copies, int count) + const VramCopy* copies, int count, bool fusedColor = false) { VkCommandBuffer cmd = VK_NULL_HANDLE; if (FAILED(g_vk.interop->BeginVkCommandBufferInterop(list, &cmd))) @@ -3561,6 +3720,27 @@ static bool record_vram_inputs(Feature& feature, ID3D12GraphicsCommandList* list const VramCopy& copy = copies[index]; VkImageLayout layout = VK_IMAGE_LAYOUT_GENERAL; g_vk.interop->GetVulkanImageLayout(copy.resource, D3D12_RESOURCE_STATE_COPY_SOURCE, &layout); + if (index == 0 && fusedColor) + { + const VkMemoryBarrier writable = {VK_STRUCTURE_TYPE_MEMORY_BARRIER, nullptr, + VK_ACCESS_MEMORY_READ_BIT | VK_ACCESS_MEMORY_WRITE_BIT, VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_SHADER_WRITE_BIT}; + g_vk.barrier(cmd, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, + 1, &writable, 0, nullptr, 0, nullptr); + VkImageMemoryBarrier sampled = {VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER}; + sampled.srcAccessMask = VK_ACCESS_MEMORY_WRITE_BIT; sampled.dstAccessMask = VK_ACCESS_SHADER_READ_BIT; + sampled.oldLayout = layout; sampled.newLayout = VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL; + sampled.srcQueueFamilyIndex = sampled.dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED; + sampled.image = copy.image; sampled.subresourceRange = {VK_IMAGE_ASPECT_COLOR_BIT, 0, 1, 0, 1}; + g_vk.barrier(cmd, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + 0, 0, nullptr, 0, nullptr, 1, &sampled); + auto& stage = feature.fusedColor[&slot - feature.inputs]; + g_fusedColorStage.record(cmd, stage.set, copy.width, copy.height, 5, slot.host[0].rowBytes / 8); + sampled.srcAccessMask = VK_ACCESS_SHADER_READ_BIT; sampled.dstAccessMask = VK_ACCESS_TRANSFER_READ_BIT; + sampled.oldLayout = VK_IMAGE_LAYOUT_SHADER_READ_ONLY_OPTIMAL; sampled.newLayout = layout; + g_vk.barrier(cmd, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, + 0, 0, nullptr, 0, nullptr, 1, &sampled); + continue; + } VkBufferImageCopy region = full_region(copy); static const Plane kPlanes[4] = {Plane::Color, Plane::Depth, Plane::Motion, Plane::Exposure}; if (feature.linearInputs) @@ -3604,7 +3784,11 @@ static bool record_vram_inputs(Feature& feature, ID3D12GraphicsCommandList* list g_vk.copyImageToBuffer(cmd, conversion, VK_IMAGE_LAYOUT_TRANSFER_SRC_OPTIMAL, slot.vram[index].buffer, 1, ®ion); } - g_vk.barrier(cmd, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, 0, 1, &kAfterTransfer, 0, + const VkMemoryBarrier completed = {VK_STRUCTURE_TYPE_MEMORY_BARRIER, nullptr, + VK_ACCESS_TRANSFER_WRITE_BIT | VK_ACCESS_SHADER_WRITE_BIT, VK_ACCESS_MEMORY_READ_BIT | VK_ACCESS_MEMORY_WRITE_BIT}; + g_vk.barrier(cmd, fusedColor ? VK_PIPELINE_STAGE_TRANSFER_BIT | VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT : + VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, 0, 1, + fusedColor ? &completed : &kAfterTransfer, 0, nullptr, 0, nullptr); return SUCCEEDED(g_vk.interop->EndVkCommandBufferInterop(list)); } @@ -3671,6 +3855,54 @@ static bool record_vram_output(ID3D12GraphicsCommandList* list, const VramBuffer static VkResult signal_split_semaphore(Feature& feature, uint64_t value); +// Only the private wait word is released; never write the game's original readback marker. No HIP call +// belongs here: HIP may hold its queue lock while NGX submits, and this runs independently of that worker. +// The coherent allocation stays alive through CUDA drain and through recorded command-list retirement. +static void release_ngx_input_wait_locked(Feature& feature, const char* reason, bool disable) +{ + if (!feature.ngxWaitActive || feature.ngxWaitCancelled) + return; + feature.ngxWaitCancelled = true; + if (feature.ngxTimeline != nullptr) g.cu.timelineCancel(feature.ngxTimeline); + if (disable) + feature.ngxGpuWait.store(false); + auto* word = static_cast(feature.ngxInputSync.mapped) + 1 + feature.ngxWaitSlot; + *word = feature.ngxWaitFrame; + MemoryBarrier(); + logf("NGX frame %u: %s; CPU released the coherent GPU-side input wait; dropping after stream drain", + feature.ngxWaitFrame, reason); +} + +static void ngx_input_watchdog() +{ + for (;;) + { + d4r_sleep_us(1000); + // Registry ownership prevents feature destruction while we inspect or release its wait. The + // worker never holds ngxWaitMutex across a CUDA call, including the wait enqueue itself. + std::lock_guard registryLock(g_featuresMutex); + for (Feature* feature : g_features) + { + std::lock_guard lock(feature->ngxWaitMutex); + if (!feature->ngxWaitActive || feature->ngxWaitCancelled) + continue; + const double elapsedUs = std::chrono::duration( + ProfileClock::now() - feature->ngxWaitStart).count(); + const uint32_t aggregate = feature->markerValue[0]; + MemoryBarrier(); + // Read the exact slot after the aggregate, as the CPU path's recheck does: a marker that + // landed between the reads must not be mistaken for a discarded command list. + const auto reason = d4r_input_wait_abort(feature->ngxWaitFrame, + feature->markerValue[1 + feature->ngxWaitSlot], aggregate, feature->retiring.load(), elapsedUs); + if (reason != D4rInputWaitAbort::None) + release_ngx_input_wait_locked(*feature, + reason == D4rInputWaitAbort::Retired ? "the feature retired" : + reason == D4rInputWaitAbort::Discarded ? "its command list was discarded" : + "the input wait or evaluation timed out (5 s)", reason == D4rInputWaitAbort::Timeout); + } + } +} + static void release_split_frame(Feature* feature, uint32_t frame) { std::lock_guard lock(feature->splitMutex); @@ -3773,6 +4005,20 @@ static bool ensure_conversion_image(Feature& feature, VramImage& image, UINT wid static void release_vram(Feature& feature) { + for (auto& stage : feature.fusedColor) feature.retiredFusedColor.push_back(stage); + for (auto& stage : feature.retiredFusedColor) + { + g_fusedColorStage.free(stage.set); + if (stage.view != VK_NULL_HANDLE) g_vk.destroyView(g_vk.device, stage.view, nullptr); + if (stage.owner != nullptr) stage.owner->Release(); + } + feature.retiredFusedColor.clear(); + if (feature.ngxTimeline != nullptr) + { + g.cu.timelineDestroy(feature.ngxTimeline); + feature.ngxTimeline = nullptr; + } + recycle_vram_buffer(feature.ngxInputSync); if (feature.splitSemaphore != VK_NULL_HANDLE) g_vk.destroySemaphore(g_vk.device, feature.splitSemaphore, nullptr); for (InputSlot& slot : feature.inputs) @@ -5645,10 +5891,21 @@ static bool wait_for_release(const std::atomic& flag) } static void release_split_frame(Feature* feature, uint32_t frame); +static bool complete_ngx_input_wait(Feature& feature, uint32_t frame, bool abort); // Every queued frame ends in exactly one call of this (published or dropped). static void frame_retired(Feature* feature, uint32_t frame) { + // Covers every early error exit, even failures after the wait but before an output event was recorded. + // Never release the split timeline or input ownership while queued NGX writes can still be running. + complete_ngx_input_wait(*feature, frame, true); + if (feature->gpuCompletion) + { + // Every queued frame ends here after its GPU work completed or was drained: release a game queue that may + // be polling for it (success paths wrote the same value from the GPU already). + static_cast(feature->ngxInputSync.mapped)[kCompletionWord + frame % kOutputSlots] = frame; + MemoryBarrier(); + } if (feature->split) release_split_frame(feature, frame); // The denoiser inputs this frame bound belong to the slot it used, and the caller stops @@ -5944,12 +6201,12 @@ static CudaObject prepare_dilated_motion(Feature& feature, InputSlot& slot, cons // into the CUDA arrays on the GPU. With D4R_SHIM_VRAM_VERIFY also compares // them with the host-staged planes of the same frame. static bool upload_inputs_vram(Feature& feature, InputSlot& slot, bool hasExposure, uint32_t frame, - FrameTiming* timing) + FrameTiming* timing, bool gpuWait = false) { const Plane planes[4] = {Plane::Color, Plane::Depth, Plane::Motion, Plane::Exposure}; CudaImage* images[4] = {&feature.color, &feature.depth, &feature.motion, &feature.exposure}; const int count = hasExposure ? 4 : 3; - const bool verify = env_uint("D4R_SHIM_VRAM_VERIFY", 0) != 0; + const bool verify = !gpuWait && env_uint("D4R_SHIM_VRAM_VERIFY", 0) != 0; const auto start = ProfileClock::now(); int result = 0; for (int index = 0; index < count && result == 0; ++index) @@ -5986,11 +6243,13 @@ static bool upload_inputs_vram(Feature& feature, InputSlot& slot, bool hasExposu copy.dstArray = images[index]->array; copy.WidthInBytes = geometry.rowBytes; copy.Height = geometry.height; - result = copy_2d(copy); + // An active GPU wait requires the known null-stream async route. Do not silently fall back to + // a synchronous or differently ordered copy while this frame's input boundary is outstanding. + result = gpuWait ? g.cu.memcpy2DAsync(©, nullptr) : copy_2d(copy); // D4R_SHIM_VRAM_ARRAY_VERIFY=1: read the CUDA array DLSS samples back // and compare it with the imported buffer it was filled from. static const bool arrayVerify = env_uint("D4R_SHIM_VRAM_ARRAY_VERIFY", 0) != 0; - if (arrayVerify && result == 0 && g.cu.ctxSynchronize() == 0) + if (!gpuWait && arrayVerify && result == 0 && g.cu.ctxSynchronize() == 0) { std::vector fromBuffer(geometry.size()), fromArray(geometry.size()); CudaMemcpy2D readBuffer = {}; @@ -6028,7 +6287,12 @@ static bool upload_inputs_vram(Feature& feature, InputSlot& slot, bool hasExposu if (timing != nullptr) timing->h2dTotal = profile_ms(start, ProfileClock::now()); if (result != 0) + { logf("VRAM input copy failed: %d", result); + if (gpuWait && feature.ngxGpuWait.exchange(false)) + logf("NGX feature %u: null-stream async input copy failed; using CPU marker polling from now on", + feature.handle.Id); + } return result == 0; } @@ -6104,7 +6368,7 @@ static int synchronize_default_stream(Feature& feature, bool* querySleep, double // A frame recorded for packed output whose output kernel then did not pack it (it was not the native one // after all): the RGBA16F result is converted on the host, rounding as the kernel store and the game-side // blit do. Slow, and only expected around a kernel change. -static int pack_output_on_host(Feature& feature, CudaMemcpy2D copy, size_t packedPitch, uint32_t frame) +static int pack_output_on_host(Feature& feature, CudaMemcpy2D copy, size_t packedPitch, uint32_t frame, bool rgba8 = false) { const size_t width = feature.output.width, height = feature.output.height; std::vector halves(width * height * 4); @@ -6123,8 +6387,11 @@ static int pack_output_on_host(Feature& feature, CudaMemcpy2D copy, size_t packe return static_cast(value * scale + 0.5f); }; for (size_t texel = 0; texel < packed.size(); ++texel) - packed[texel] = unorm(halves[texel * 4], 1023.0f) | unorm(halves[texel * 4 + 1], 1023.0f) << 10 | - unorm(halves[texel * 4 + 2], 1023.0f) << 20 | unorm(halves[texel * 4 + 3], 3.0f) << 30; + packed[texel] = rgba8 + ? unorm(halves[texel * 4], 255.0f) | unorm(halves[texel * 4 + 1], 255.0f) << 8 | + unorm(halves[texel * 4 + 2], 255.0f) << 16 | unorm(halves[texel * 4 + 3], 255.0f) << 24 + : unorm(halves[texel * 4], 1023.0f) | unorm(halves[texel * 4 + 1], 1023.0f) << 10 | + unorm(halves[texel * 4 + 2], 1023.0f) << 20 | unorm(halves[texel * 4 + 3], 3.0f) << 30; CudaMemcpy2D upload = {}; upload.srcMemoryType = CUDA_MEMORY_HOST; upload.srcHost = packed.data(); @@ -6139,6 +6406,33 @@ static int pack_output_on_host(Feature& feature, CudaMemcpy2D copy, size_t packe return result; } +// Worker only: finish the tracked wait once output completed, or release and drain an aborted submission. +static bool complete_ngx_input_wait(Feature& feature, uint32_t frame, bool abort) +{ + { + std::lock_guard lock(feature.ngxWaitMutex); + if (!feature.ngxWaitActive || feature.ngxWaitFrame != frame) + return true; + if (abort) + release_ngx_input_wait_locked(feature, "submission or output completion failed", false); + } + // Success callers already synchronized the output; error callers may have no completion event, so + // drain the real context (the bridge export bypasses NGX sync elision) after releasing the wait. Keep + // the watchdog armed during this call, and the slot/resources owned until it returns. + if (abort && g.cu.ctxSynchronize() != 0) + { + feature.ngxGpuWait.store(false); + feature.retiring.store(true); + logf("NGX frame %u: context drain failed after input-wait release; retiring feature", frame); + } + std::lock_guard lock(feature.ngxWaitMutex); + const bool valid = !feature.ngxWaitCancelled; + feature.ngxWaitActive = false; + if (!valid) + feature.evaluatedFrames = 0; // a discarded evaluation must not supply temporal history to the next one + return valid; +} + // VRAM interop, worker thread: copies the finished result into a free output // slot's imported buffer and publishes it; evaluate() copies it on to the game. static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& params, FrameTiming& timing, @@ -6174,13 +6468,57 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa int result = rowBytes / (params.packedOutput ? 2 : 1) * feature->output.height > buffer.bytes ? -1 : feature->outputRedirected ? 0 : params.packedOutput ? pack_output_on_host(*feature, copy, params.outputInPlace != 0 ? copy.dstPitch - : rowBytes / 2, frame) + : rowBytes / 2, frame, params.packedRgba8) : copy_2d(copy); + // GPU-side completion: the word is written behind every output write on NGX's stream. A failed enqueue is + // covered by frame_retired's CPU write once this frame has synchronized below. + if (result == 0 && params.gpuCompletion && !feature->completionFired && + g.cu.ngxWriteValue32(feature->ngxInputSync.device + sizeof(uint32_t) * (kCompletionWord + frame % kOutputSlots), + frame) != 0) + { + static bool logged = false; + if (!logged) + logf("frame %u: GPU completion write could not be queued; the CPU releases the game's queue", frame); + logged = true; + } + bool nativeQueued = false; + if (result == 0 && params.split && feature->ngxNativeCompletion && !feature->retiring.load()) + { + bool head; + { + std::lock_guard lock(feature->splitMutex); + head = feature->splitCompletion.can_signal(frame); + } + // The callback is ordered after every output write. It releases this frame on the GPU's + // completion boundary and wakes the native condition; the worker retains resources until + // that condition completes. Future discarded frames cannot cross an earlier producer. + if (head) + { + nativeQueued = g.cu.timelineQueue(feature->ngxTimeline, frame) == 0; + if (!nativeQueued) + { + feature->ngxNativeCompletion = false; + logf("NGX native completion enqueue failed; using CPU output completion from now on"); + } + } + } if (result == 0) { const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; const double cpuStart = timing.enabled ? profile_thread_cpu_ms() : -1.0; - result = synchronize_default_stream(*feature, &timing.outputSyncBlocking, &feature->publishWaitEmaUs); + if (nativeQueued) + { + result = g.cu.timelineWait(feature->ngxTimeline); + timing.outputSyncNative = result == 0; + if (result != 0) + { + g.cu.timelineCancel(feature->ngxTimeline); + feature->ngxNativeCompletion = false; + logf("NGX native completion wait failed (%d); using CPU output completion from now on", result); + } + } + if (!nativeQueued || result != 0) + result = synchronize_default_stream(*feature, &timing.outputSyncBlocking, &feature->publishWaitEmaUs); if (timing.enabled) { timing.outputSyncWall = profile_ms(syncStart, ProfileClock::now()); @@ -6189,11 +6527,29 @@ static void publish_vram(Feature* feature, uint32_t frame, const FrameParams& pa timing.outputSyncCpu = std::max(0.0, cpuEnd - cpuStart); } } + if (result == 0 && !complete_ngx_input_wait(*feature, frame, false)) + { + frame_retired(feature, frame); + return; // queued work drained, but watchdog-released input is not a publishable frame + } if (result == 0 && timing.gpuEventsRecorded) { float gpuEvalMs = -1.0f; if (g.cu.eventElapsedTime(&gpuEvalMs, feature->profileStart, feature->profileEnd) == 0) timing.gpuEval = gpuEvalMs; + float dlssMs = -1.0f, inputMs = -1.0f; + if (timing.inputEventRecorded && + g.cu.eventElapsedTime(&dlssMs, feature->profileInputsReleased, feature->profileEnd) == 0 && + g.cu.eventElapsedTime(&inputMs, feature->profileStart, feature->profileInputsReleased) == 0) + { + timing.dlssGpu = dlssMs; + timing.inputGpuWait = inputMs; + } + else if (!timing.gpuInputWait && timing.gpuEval >= 0.0) + { + timing.dlssGpu = timing.gpuEval; + timing.inputGpuWait = 0.0; + } } if (result != 0) { @@ -6316,6 +6672,18 @@ static bool recreate_for_render_size(Feature& feature, FrameParams& params, uint static bool wait_for_input_marker(Feature* feature, int slotIndex, uint32_t frame); +static bool ngx_gpu_wait_eligible(Feature& feature, const FrameParams& params, uint32_t frame) +{ + char directory[MAX_PATH]; + // The denoiser can mix GPU guides with CPU-transcoded auxiliary guides and alpha readbacks. Keep + // that mixed route on the CPU readiness contract; never infer that main-plane VRAM means all guides + // are GPU-only. Likewise verification, captures, dilation and host staging must see real input first. + return feature.ngxGpuWait.load() && params.gpuInputMarker && params.vram && params.split && !feature.rayReconstruction && + env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0 && env_uint("D4R_SHIM_VRAM_ARRAY_VERIFY", 0) == 0 && + env_uint("D4R_MOTION_DILATION", 0) == 0 && !input_dump_directory(frame, directory) && + (feature.linearInputs || (env_uint("D4R_SHIM_ASYNC_2D", 1) != 0 && g.cu.memcpy2DAsync != nullptr)); +} + static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, FrameParams params, FrameTiming timing) { InputSlot& slot = feature->inputs[slotIndex]; @@ -6324,19 +6692,102 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram timing.workerWait = profile_ms(timing.workerQueued, start); FrameTiming* stages = timing.enabled ? &timing : nullptr; const bool overlapped = async_copies() && copy_streams() != nullptr; + if (feature->ngxNativeCompletion && g.cu.timelinePrepare(feature->ngxTimeline) != 0) + { + feature->ngxNativeCompletion = false; + logf("NGX native completion still busy; using CPU output completion from now on"); + } + int eventStartResult = -1; if (timing.markerDeferred) { const auto markerStart = ProfileClock::now(); - if (!wait_for_input_marker(feature, slotIndex, frame)) - return; - if (timing.enabled) - timing.markerWait = profile_ms(markerStart, ProfileClock::now()); + if (ngx_gpu_wait_eligible(*feature, params, frame)) + { + complete_pending_download(); // earlier host staging must not synchronize against our new wait + // Mint textures/arrays before blocking the null stream. Lazy allocation can synchronize + // internally; it must not be part of the steady-state input-readiness submission window. + bool ready = feature->linearInputs ? ensure_linear_inputs(slot, params.hasExposure) : true; + if (!feature->linearInputs) + { + const Plane planes[] = {Plane::Color, Plane::Depth, Plane::Motion, Plane::Exposure}; + CudaImage* images[] = {&feature->color, &feature->depth, &feature->motion, &feature->exposure}; + for (int index = 0; index < (params.hasExposure ? 4 : 3) && ready; ++index) + ready = ensure_cuda_image(*images[index], slot.host[index].width, slot.host[index].height, + plane_format(planes[index]), plane_channels(planes[index]), false, index == 0); + } + ready = ready && ensure_cuda_image(feature->output, feature->outputWidth, feature->outputHeight, + CUDA_FORMAT_HALF, 4, true, false); + if (!ready) + { + slot.busy = false; + frame_retired(feature, frame); + return; + } + if (!feature->ngxWaitEventsWarmed) + { + if (feature->outputReadyEvent != nullptr) + g.cu.eventRecord(feature->outputReadyEvent, nullptr); + if (timing.enabled && feature->profileEventsReady) + g.cu.eventRecord(feature->profileEnd, nullptr); + feature->ngxWaitEventsWarmed = true; + } + eventStartResult = timing.enabled && feature->profileEventsReady + ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; + { + std::lock_guard lock(feature->ngxWaitMutex); + feature->ngxWaitActive = true; + feature->ngxWaitCancelled = false; + feature->ngxWaitFrame = frame; + feature->ngxWaitSlot = slotIndex; + feature->ngxWaitStart = ProfileClock::now(); + } + // Arm before the call: a blocked HIP enqueue/NGX evaluate must never prevent timeout release. + const int queued = g.cu.streamWaitValue32Eq( + feature->ngxInputSync.device + sizeof(uint32_t) * (1 + slotIndex), frame); + if (queued == 0) + { + timing.gpuInputWait = true; + // Profiling only: a timing event behind the wait completes when the game's input word is released, + // so it splits gpu_eval into input readiness and DLSS's own work. Recording queues one marker on + // the null stream and never waits on the CPU; elapsed times are read after output completion. + timing.inputEventRecorded = eventStartResult == 0 && feature->profileInputsReleased != nullptr && + g.cu.eventRecord(feature->profileInputsReleased, nullptr) == 0; + bool cancelled; + { + std::lock_guard lock(feature->ngxWaitMutex); + cancelled = feature->ngxWaitCancelled; + } + if (cancelled) + { + complete_ngx_input_wait(*feature, frame, true); + slot.busy = false; + frame_retired(feature, frame); + return; + } + } + else + { + // Release and drain even on rejection: a partially submitted driver operation must + // never survive a switch back to CPU polling. Do not hold the ticket lock over HIP. + complete_ngx_input_wait(*feature, frame, true); + feature->ngxGpuWait.store(false); + logf("NGX feature %u: GPU-side input wait unavailable (%d); using CPU marker polling from now on", + feature->handle.Id, queued); + } + } + if (!timing.gpuInputWait) + { + if (!wait_for_input_marker(feature, slotIndex, frame)) + return; + if (timing.enabled) + timing.markerWait = profile_ms(markerStart, ProfileClock::now()); + } } - if (params.vram && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) + if (!timing.gpuInputWait && params.vram && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) dump_vram_inputs(slot, frame, params, params.hasExposure ? 4 : 3); const bool linearInputs = params.vram && feature->linearInputs; bool ok = linearInputs ? ensure_linear_inputs(slot, params.hasExposure) - : params.vram ? upload_inputs_vram(*feature, slot, params.hasExposure, frame, stages) + : params.vram ? upload_inputs_vram(*feature, slot, params.hasExposure, frame, stages, timing.gpuInputWait) : overlapped ? upload_inputs_async(*feature, slot, params.hasExposure, frame, stages) : upload_plane(Plane::Color, slot.host[0], feature->color, frame, stages, 0) && @@ -6347,7 +6798,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram ok = ok && ensure_cuda_image(feature->output, feature->outputWidth, feature->outputHeight, CUDA_FORMAT_HALF, 4, true, false); const auto dilationStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; - const CudaObject dilatedMotion = ok ? prepare_dilated_motion(*feature, slot, params) : 0; + const CudaObject dilatedMotion = ok && !timing.gpuInputWait ? prepare_dilated_motion(*feature, slot, params) : 0; // Ray Reconstruction: the auxiliary surfaces are copied in the same window as the four // main planes, so the one synchronize below covers them as well. The lock is taken only // around the copies and, later, around the evaluate - never across the synchronize, which @@ -6363,19 +6814,21 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram } if (!auxUploaded) { + complete_ngx_input_wait(*feature, frame, true); slot.busy = false; frame_retired(feature, frame); return; } } // Device-to-array copies (VRAM interop, staged uploads) return before they - // finish, and NGX does not evaluate on the stream they were issued on, so - // without this DLSS can sample a mix of this frame's and the previous + // finish. Staged copies can use another stream, so without their completion + // DLSS can sample a mix of this frame's and the previous // frame's colour/depth/motion: invisible when the camera is still, beaded // thin geometry in motion. D4R_SHIM_INPUT_SYNC=0 restores the old behaviour. static const bool inputSync = env_uint("D4R_SHIM_INPUT_SYNC", 1) != 0; - // Linear inputs are already ready; other copies need completion before NGX samples them. - if (ok && (dilatedMotion != 0 || auxUploaded || (inputSync && !linearInputs))) + // Linear inputs are ready after the marker. The GPU-wait route orders all copies and NGX on + // the null stream itself, so its input-sync CPU wait would only put readiness back on the critical path. + if (ok && !timing.gpuInputWait && (dilatedMotion != 0 || auxUploaded || (inputSync && !linearInputs))) { const int syncResult = synchronize_default_stream(*feature); if (syncResult != 0) @@ -6395,6 +6848,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram complete_pending_download(); if (!ok) { + complete_ngx_input_wait(*feature, frame, true); slot.busy = false; frame_retired(feature, frame); return; @@ -6415,6 +6869,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram if (!publish_alpha_output(*feature, slot)) { auxLock.unlock(); + complete_ngx_input_wait(*feature, frame, true); slot.busy = false; frame_retired(feature, frame); return; @@ -6450,7 +6905,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram else if (params.vram && params.split && feature->outputs[frame % kOutputSlots].vram.bytes >= bytes) target = feature->outputs[frame % kOutputSlots].vram.device; if (params.packedOutput) - rowBytes |= 1; // the bridge's flag for R10G10B10A2 texels (row pitches are multiples of 8) + rowBytes |= params.packedRgba8 ? 2 : 1; // R4: RGBA8, R3: RGB10A2; row pitches are multiples of 8 else if (!params.vram && overlapped && ensure_linear(feature->outputLinear, feature->outputLinearBytes, bytes)) target = feature->outputLinear; const bool redirect = nativeOutput && target != 0 && @@ -6465,8 +6920,14 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram } feature->outputRedirected = redirect; } - const int eventStartResult = timing.enabled && feature->profileEventsReady - ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; + if (!timing.gpuInputWait) + eventStartResult = timing.enabled && feature->profileEventsReady + ? g.cu.eventRecord(feature->profileStart, nullptr) : -1; + // GPU completion: with a redirected output the native output kernel is the last writer the game reads, so + // the bridge may queue the completion word right behind it (it checks the kernel actually launched). + if (params.gpuCompletion && feature->outputRedirected && g.cu.ngxArmCompletion != nullptr) + g.cu.ngxArmCompletion(feature->ngxInputSync.device + sizeof(uint32_t) * (kCompletionWord + frame % kOutputSlots), + frame, params.packedRgba8 ? 1 : 0); const auto evalStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; // The core dispatches feature 13 to the denoiser it loaded, so this call is the same // either way; only the feature id that created it differs. @@ -6475,13 +6936,16 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram result = g.ngx.evaluateFeature(feature->cudaHandle, feature->cudaParams, nullptr); if (auxLock.owns_lock()) auxLock.unlock(); - if (feature->outputRedirected && g.cu.outputKernelNative != nullptr && g.cu.outputKernelNative() != 1) + feature->completionFired = g.cu.ngxCompletionFired != nullptr && g.cu.ngxCompletionFired() == 1; + if (feature->outputRedirected && ((g.cu.outputKernelNative != nullptr && g.cu.outputKernelNative() != 1) || + (params.packedRgba8 && (g.cu.outputKernelRgba8 == nullptr || g.cu.outputKernelRgba8() != 1)))) { // this frame's output kernel ignored the redirect: publish from the array as usual g.cu.setArrayRedirect(feature->output.array, 0, 0); feature->outputRedirected = false; } feature->packedReady = feature->outputRedirected; + feature->rgba8Ready = feature->outputRedirected && g.cu.outputKernelRgba8 != nullptr && g.cu.outputKernelRgba8() == 1; const auto evalReturned = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; const int eventEndResult = eventStartResult == 0 ? g.cu.eventRecord(feature->profileEnd, nullptr) : -1; const auto syncStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; @@ -6499,6 +6963,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram if (result != NGX_SUCCESS || syncResult != 0) { logf("frame %u: CUDA_EvaluateFeature -> 0x%08x, sync %d", frame, result, syncResult); + complete_ngx_input_wait(*feature, frame, true); slot.busy = false; frame_retired(feature, frame); return; @@ -6508,6 +6973,7 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram // alpha would leave the caller believing it was written. if (feature->rayReconstruction && !download_alpha_output(*feature, frame, params)) { + complete_ngx_input_wait(*feature, frame, true); slot.busy = false; frame_retired(feature, frame); return; @@ -6517,6 +6983,11 @@ static void run_evaluation(Feature* feature, int slotIndex, uint32_t frame, Fram timing.ngxHost = profile_ms(evalStart, evalReturned); timing.ctxSync = profile_ms(syncStart, evaluated); timing.gpuEval = gpuEvalMs; + if (gpuEvalMs >= 0.0f && !timing.gpuInputWait) + { + timing.dlssGpu = gpuEvalMs; + timing.inputGpuWait = 0.0; + } timing.gpuEventsRecorded = gpuEventsRecorded && params.vram && !eval_sync_enabled(); } if (params.vram) @@ -6566,17 +7037,21 @@ static bool wait_for_input_marker(Feature* feature, int slotIndex, uint32_t fram { InputSlot& slot = feature->inputs[slotIndex]; const auto start = std::chrono::steady_clock::now(); + static const bool fixedOverride = GetEnvironmentVariableA("D4R_SHIM_MARKER_POLL_US", nullptr, 0) != 0; + static const unsigned fixedPoll = env_uint("D4R_SHIM_MARKER_POLL_US", 200); + const double predictedUs = feature->markerWaitEmaUs.load(); + double lastNotReadyUs = -1.0; // A later input marker does not prove this slot's copies were submitted: // the game may have discarded this frame's command list. Require its own // tag, otherwise stale input data can enter DLSS's temporal history. const volatile uint32_t* inputMarker = feature->markerValue + 1 + slotIndex; - while (*inputMarker != frame) + while (*inputMarker != frame || feature->retiring.load()) { std::atomic_thread_fence(std::memory_order_acquire); const bool skipped = static_cast(*feature->markerValue - frame) > 0; // The GPU may have filled our slot between the loop condition and // the aggregate-marker read. Recheck before deciding it was skipped. - if (*inputMarker == frame) + if (*inputMarker == frame && !feature->retiring.load()) break; if (feature->retiring.load() || skipped || std::chrono::steady_clock::now() - start > std::chrono::seconds(5)) { @@ -6587,12 +7062,18 @@ static bool wait_for_input_marker(Feature* feature, int slotIndex, uint32_t fram frame_retired(feature, frame); return false; } - // D4R_SHIM_MARKER_POLL_US: poll interval (0 = yield); detection latency is on the critical path - static const unsigned markerPollUs = env_uint("D4R_SHIM_MARKER_POLL_US", 200); - if (markerPollUs == 0) - std::this_thread::yield(); - else - d4r_sleep_us(markerPollUs); + lastNotReadyUs = std::chrono::duration(ProfileClock::now() - start).count(); + // An explicit N retains fixed pacing (0 yields). Otherwise learn readiness per feature, sleep + // in <=200 us chunks, yield near prediction, and back off when late, as the output wait does. + d4r_sleep_us(fixedOverride ? fixedPoll + : d4r_marker_poll_delay_us(predictedUs, lastNotReadyUs)); + } + if (!fixedOverride) + { + const double waitedUs = std::chrono::duration(ProfileClock::now() - start).count(); + // Estimate readiness between the last not-ready read and this one, not the end of our sleep. + const double sample = std::min(lastNotReadyUs >= 0.0 ? 0.5 * (lastNotReadyUs + waitedUs) : waitedUs, 20000.0); + feature->markerWaitEmaUs.store(predictedUs == 0.0 ? sample : 0.8 * predictedUs + 0.2 * sample); } return true; } @@ -6606,7 +7087,8 @@ static void prepare_inputs(Feature* feature, int slotIndex, uint32_t frame, Fram // uploads and diagnostic readbacks; waiting only before NGX would copy stale data into its arrays. static const bool deferMarker = env_uint("D4R_SHIM_DEFER_MARKER", 1) != 0 && env_uint("D4R_MOTION_DILATION", 0) == 0; - if (params.vram && deferMarker && env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) + if (params.vram && (deferMarker || ngx_gpu_wait_eligible(*feature, params, frame)) && + env_uint("D4R_SHIM_VRAM_VERIFY", 0) == 0) { timing.markerDeferred = true; if (timing.enabled) @@ -6705,7 +7187,8 @@ static void finish_publish(Feature* feature, int target, uint32_t frame, const F const uint32_t completed = ++feature->completedFrames; frame_retired(feature, frame); const auto finished = ProfileClock::now(); - if (timing.enabled) + static const bool profileLog = env_uint("D4R_PROFILE_LOG", 1) != 0; // 0: time everything, write no lines + if (timing.enabled && profileLog) { static thread_local ProfileClock::time_point lastFinished; const double interval = lastFinished.time_since_epoch().count() != 0 @@ -6720,11 +7203,12 @@ static void finish_publish(Feature* feature, int target, uint32_t frame, const F "d3d_input_record=%.3f d3d_output_record=%.3f game_call=%.3f " "color_convert=%.3f depth_convert=%.3f motion_convert=%.3f exposure_convert=%.3f " "color_h2d=%.3f depth_h2d=%.3f motion_h2d=%.3f exposure_h2d=%.3f " - "ngx_host=%.3f ctx_sync=%.3f gpu_eval=%.3f output_sync_wall=%.3f " + "ngx_host=%.3f ctx_sync=%.3f gpu_eval=%.3f gpu_eval_scope=%s input_wait=%s output_sync_wall=%.3f " "output_sync_cpu=%.3f output_sync_blocking=%d output_sync_method=%s d2h=%.3f output_convert=%.3f " "h2d_total=%.3f d2h_issue=%.3f d2h_wait=%.3f " "prep=%.3f worker_wait=%.3f worker=%.3f finish_wait=%.3f finish=%.3f pipeline=%.3f " - "throughput_interval=%.3f presented=%u age=%d motion_dilation=%u motion_dilation_wall=%.3f", + "throughput_interval=%.3f presented=%u age=%d motion_dilation=%u motion_dilation_wall=%.3f " + "dlss_gpu=%.3f input_gpu_wait=%.3f vk_input=%.3f vk_dlss_wait=%.3f vk_output=%.3f vk_total=%.3f", frame, params.renderWidth, params.renderHeight, feature->outWidth, feature->outHeight, slot.host[0].width, slot.host[0].height, slot.host[1].width, slot.host[1].height, slot.host[2].width, slot.host[2].height, params.hasExposure, feature->quality, feature->flags, @@ -6734,14 +7218,18 @@ static void finish_publish(Feature* feature, int target, uint32_t frame, const F timing.inputRecord, timing.outputRecord, timing.gameCall, timing.convert[0], timing.convert[1], timing.convert[2], timing.convert[3], timing.upload[0], timing.upload[1], timing.upload[2], timing.upload[3], - timing.ngxHost, timing.ctxSync, timing.gpuEval, timing.outputSyncWall, timing.outputSyncCpu, - timing.outputSyncBlocking, timing.outputSyncBlocking ? "event_query_sleep" : "context", + timing.ngxHost, timing.ctxSync, timing.gpuEval, + timing.gpuInputWait ? "launch_to_completion" : "evaluation", timing.gpuInputWait ? "gpu" : "cpu", + timing.outputSyncWall, timing.outputSyncCpu, + timing.outputSyncBlocking, timing.outputSyncNative ? "native_callback_wait" : + timing.outputSyncBlocking ? "event_query_sleep" : "context", timing.download, timing.outputConvert, timing.h2dTotal, timing.d2hIssue, timing.d2hWait, timing.prep, timing.workerWait, timing.worker, timing.finishWait, profile_ms(start, finished), profile_ms(timing.prepStart, finished), interval, timing.presentedFrame, timing.presentedFrame != 0 ? static_cast(frame - timing.presentedFrame) : -1, - params.motionDilation, timing.motionDilationWall); + params.motionDilation, timing.motionDilationWall, timing.dlssGpu, timing.inputGpuWait, + timing.vkInput, timing.vkDlssWait, timing.vkOutput, timing.vkTotal); } if (completed <= 5 || completed % 120 == 0) logf("frame %u done (slot %d)", frame, target); @@ -7206,12 +7694,17 @@ static NgxResult create_cuda_feature(Feature& feature) else logf("output event query/sleep wait enabled for this feature"); } - if (created == NGX_SUCCESS && profile_enabled() && g.cu.eventCreate != nullptr && + // D4R_PROFILE_GPU_EVENTS=0 keeps D4R_PROFILE's CPU timings without the CUDA timing events + if (created == NGX_SUCCESS && profile_enabled() && env_uint("D4R_PROFILE_GPU_EVENTS", 1) != 0 && + g.cu.eventCreate != nullptr && g.cu.eventRecord != nullptr && g.cu.eventElapsedTime != nullptr && g.cu.eventDestroy != nullptr) { const int startResult = g.cu.eventCreate(&feature.profileStart, 0); const int endResult = startResult == 0 ? g.cu.eventCreate(&feature.profileEnd, 0) : -1; feature.profileEventsReady = startResult == 0 && endResult == 0; + // optional: without it GPU-input-wait frames report dlss_gpu=-1 and keep gpu_eval + if (feature.profileEventsReady && g.cu.eventCreate(&feature.profileInputsReleased, 0) != 0) + feature.profileInputsReleased = nullptr; if (!feature.profileEventsReady) { logf("D4R_PROFILE: CUDA event creation failed (%d, %d)", startResult, endResult); @@ -7621,6 +8114,8 @@ static VkResult signal_split_semaphore(Feature& feature, uint64_t value) { if (feature.engineNetTimeline) return static_cast(g_engineNet.signalTimeline(feature.engineNet, value)); + if (feature.ngxTimeline != nullptr) + return static_cast(g.cu.timelineSignal(feature.ngxTimeline, value)); VkSemaphoreSignalInfo info = {VK_STRUCTURE_TYPE_SEMAPHORE_SIGNAL_INFO}; info.semaphore = feature.splitSemaphore; info.value = value; @@ -7725,6 +8220,78 @@ static bool create_engine_shared_buffer(VramBuffer& target, size_t bytes, bool h return true; } +static bool initialize_ngx_gpu_wait(Feature& feature) +{ + if (env_uint("D4R_SHIM_GPU_INPUT_WAIT", 0) == 0) + return false; + // Age-zero output only; keep queue depth and the existing split-frame contract unchanged. + if (!feature.split || feature.rayReconstruction || g_vk.fill == nullptr || + g.cu.streamWaitValue32Eq == nullptr || + !create_engine_shared_buffer(feature.ngxInputSync, 256, true)) + { + logf("NGX feature %u: GPU-side input wait unavailable (needs split VRAM, Super Sampling, coherent export " + "and the matching bridge); using CPU marker polling", feature.handle.Id); + return false; + } + static std::once_flag watchdog; + std::call_once(watchdog, [] { std::thread(ngx_input_watchdog).detach(); }); + feature.ngxGpuWait.store(true); + // D4R_SHIM_GPU_COMPLETION=1 (default 0): the game's queue waits for DLSS on the GPU, polling the frame's word + // in the same coherent buffer, instead of a split list released by a CPU timeline signal (one host round trip, + // ~0.1 ms, per frame). The pool is never freed: game command lists still reference the set after retirement. + if (env_uint("D4R_SHIM_GPU_COMPLETION", 0) != 0) + { + if (feature.completionSet.set == VK_NULL_HANDLE && g_waitReady && g.cu.ngxWriteValue32 != nullptr) + feature.gpuCompletion = g_waitStage.wait_set(feature.completionSet, feature.ngxInputSync.buffer, 256); + else + feature.gpuCompletion = feature.completionSet.set != VK_NULL_HANDLE && g.cu.ngxWriteValue32 != nullptr; + for (uint32_t index = 0; index < kOutputSlots; ++index) + static_cast(feature.ngxInputSync.mapped)[kCompletionWord + index] = 0; + logf("NGX feature %u: GPU-side completion %s", feature.handle.Id, + feature.gpuCompletion ? "enabled (the game's queue polls DLSS's completion word)" + : "unavailable (wait shader or bridge export missing); using the split timeline"); + } + return true; +} + +// The game's queue holds here until DLSS wrote this frame's completion word (or the bounded poll gives up). +static bool record_gpu_completion_wait(Feature& feature, ID3D12GraphicsCommandList* list, uint32_t frame) +{ + VkCommandBuffer command = VK_NULL_HANDLE; + if (FAILED(g_vk.interop->BeginVkCommandBufferInterop(list, &command))) + return false; + g_waitStage.record_wait(command, feature.completionSet, kCompletionWord + frame % kOutputSlots, frame, + kCompletionPollLimit); + // Everything after the wait, including the output copy, sees DLSS's writes. + const VkMemoryBarrier after = {VK_STRUCTURE_TYPE_MEMORY_BARRIER, nullptr, VK_ACCESS_MEMORY_WRITE_BIT, + VK_ACCESS_MEMORY_READ_BIT | VK_ACCESS_MEMORY_WRITE_BIT}; + g_vk.barrier(command, VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, 0, 1, &after, + 0, nullptr, 0, nullptr); + return SUCCEEDED(g_vk.interop->EndVkCommandBufferInterop(list)); +} + +static bool record_ngx_input_ready(Feature& feature, ID3D12GraphicsCommandList* list, int slotIndex, uint32_t frame) +{ + if (!feature.ngxGpuWait.load()) + return false; + VkCommandBuffer command = VK_NULL_HANDLE; + if (FAILED(g_vk.interop->BeginVkCommandBufferInterop(list, &command))) + return false; + // All input copies/conversions precede this release marker. The imported inputs and the shared word + // use the same memory barrier protocol as the engine's GPU handoff. Keep this before the list split: + // placing the fill in its second half would deadlock that half against its own NGX input wait. + const VkMemoryBarrier before = {VK_STRUCTURE_TYPE_MEMORY_BARRIER, nullptr, + VK_ACCESS_MEMORY_WRITE_BIT, VK_ACCESS_TRANSFER_WRITE_BIT}; + g_vk.barrier(command, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, 1, &before, + 0, nullptr, 0, nullptr); + g_vk.fill(command, feature.ngxInputSync.buffer, sizeof(uint32_t) * (1 + slotIndex), sizeof(uint32_t), frame); + const VkMemoryBarrier after = {VK_STRUCTURE_TYPE_MEMORY_BARRIER, nullptr, VK_ACCESS_TRANSFER_WRITE_BIT, + VK_ACCESS_MEMORY_READ_BIT | VK_ACCESS_MEMORY_WRITE_BIT}; + g_vk.barrier(command, VK_PIPELINE_STAGE_TRANSFER_BIT, VK_PIPELINE_STAGE_ALL_COMMANDS_BIT, 0, 1, &after, + 0, nullptr, 0, nullptr); + return SUCCEEDED(g_vk.interop->EndVkCommandBufferInterop(list)); +} + // Only once no command list uses the engine and no network launch is pending. static void release_engine_network(Feature& feature) { @@ -8839,6 +9406,74 @@ static bool engine_evaluate(Feature& feature, ID3D12GraphicsCommandList* list, I return true; } +static bool profile_timestamps_ready(Feature& feature) +{ + if (feature.timestampHeap != nullptr) + return true; + if (feature.timestampFailed || g.device == nullptr) + return false; + feature.timestampFailed = true; + D3D12_COMMAND_QUEUE_DESC queueDesc = {}; + queueDesc.Type = D3D12_COMMAND_LIST_TYPE_DIRECT; + ID3D12CommandQueue* queue = nullptr; + UINT64 frequency = 0; + if (SUCCEEDED(g.device->CreateCommandQueue(&queueDesc, __uuidof(ID3D12CommandQueue), reinterpret_cast(&queue)))) + { + queue->GetTimestampFrequency(&frequency); + queue->Release(); + } + D3D12_QUERY_HEAP_DESC heapDesc = {}; + heapDesc.Type = D3D12_QUERY_HEAP_TYPE_TIMESTAMP; + heapDesc.Count = 4 * kOutputSlots; + uint8_t* mapped = nullptr; + if (frequency == 0 || + FAILED(g.device->CreateQueryHeap(&heapDesc, __uuidof(ID3D12QueryHeap), reinterpret_cast(&feature.timestampHeap)))) + { + logf("D4R_PROFILE: game-queue timestamps unavailable (frequency %llu)", static_cast(frequency)); + feature.timestampHeap = nullptr; + return false; + } + if (!create_buffer(D3D12_HEAP_TYPE_READBACK, 8 * 4 * kOutputSlots, &feature.timestampReadback, &mapped)) + { + feature.timestampHeap->Release(); + feature.timestampHeap = nullptr; + return false; + } + feature.timestampValues = reinterpret_cast(mapped); + for (int index = 0; index < 4 * kOutputSlots; ++index) + feature.timestampValues[index] = 0; + feature.timestampTicksPerMs = static_cast(frequency) / 1000.0; + feature.timestampFailed = false; + logf("D4R_PROFILE: game-queue timestamps enabled (%llu Hz)", static_cast(frequency)); + return true; +} + +static void profile_timestamp(Feature& feature, ID3D12GraphicsCommandList* list, uint32_t frame, int point) +{ + const UINT base = 4 * (frame % kOutputSlots); + list->EndQuery(feature.timestampHeap, D3D12_QUERY_TYPE_TIMESTAMP, base + point); + if (point == 3) + list->ResolveQueryData(feature.timestampHeap, D3D12_QUERY_TYPE_TIMESTAMP, base, 4, feature.timestampReadback, + 8ull * base); +} + +// The evaluation kOutputSlots frames earlier, if its resolve has landed; clears the slot for this frame. +static void read_profile_timestamps(Feature& feature, uint32_t frame, FrameTiming& timing) +{ + volatile uint64_t* t = feature.timestampValues + 4 * (frame % kOutputSlots); + const uint64_t t0 = t[0], t1 = t[1], t2 = t[2], t3 = t[3]; + if (t0 != 0 && t0 <= t1 && t1 <= t2 && t2 <= t3) + { + const double perMs = feature.timestampTicksPerMs; + timing.vkInput = (t1 - t0) / perMs; + timing.vkDlssWait = (t2 - t1) / perMs; + timing.vkOutput = (t3 - t2) / perMs; + timing.vkTotal = (t3 - t0) / perMs; + } + for (int index = 0; index < 4; ++index) + t[index] = 0; +} + static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* handle, void* parameters) { std::lock_guard apiLock(g_featureApiMutex); @@ -9060,9 +9695,23 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand g.cu.registerLinearTexture != nullptr; if (feature->linearInputs) logf("linear inputs: NGX samples the interop buffers directly"); + const bool gpuWait = p.vram && initialize_ngx_gpu_wait(*feature); + if (feature->split && !feature->rayReconstruction && !feature->engineNetTimeline && + env_uint("D4R_SHIM_NATIVE_COMPLETION", 0) != 0 && g.cu.timelineCreate != nullptr && + g.cu.timelinePrepare != nullptr && g.cu.timelineCancel != nullptr && g.cu.timelineQueue != nullptr && + g.cu.timelineSignal != nullptr && g.cu.timelineDestroy != nullptr && g.cu.timelineWait != nullptr) + { + feature->ngxNativeCompletion = g.worker.call([&] { + return g.cu.timelineCreate(g_vk.device, reinterpret_cast(feature->splitSemaphore), + &feature->ngxTimeline); + }) == 0; + logf("NGX output completion for this feature: %s", feature->ngxNativeCompletion + ? "native HIP callback signals the split timeline" : "CPU event polling (native callback unavailable)"); + } logf("VRAM interop %s for this feature%s%s", p.vram ? "on" : "off", feature->split ? ", presenting each frame's own result (split frames)" : "", - p.vram ? ", waiting for input readiness with CPU sleeps" : ""); + gpuWait ? ", GPU-side input wait on NGX's null stream (CPU waits for diagnostics/host input reads)" : + p.vram ? ", waiting for input readiness with adaptive CPU polling (or explicit MarkerPollUs)" : ""); } else if (feature->vram && !p.vram) { @@ -9096,11 +9745,17 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand outputDesc.Format == DXGI_FORMAT_R10G10B10A2_UNORM && vramOutput.width % 2 == 0 && env_uint("D4R_SHIM_OUTPUT_DIRECT", 0) != 0 && env_uint("D4R_SHIM_OUTPUT_PACKED", 0) != 0 && feature->packedReady.load(); + p.packedRgba8 = feature->split && feature->outputDirectAllowed && feature->preset == 11 && + outputDesc.Format == DXGI_FORMAT_R8G8B8A8_UNORM && vramOutput.width % 2 == 0 && + env_uint("D4R_SHIM_OUTPUT_DIRECT", 0) != 0 && env_uint("D4R_SHIM_OUTPUT_RGBA8", 0) != 0 && + feature->rgba8Ready.load(); + p.packedOutput |= p.packedRgba8; static bool packedLogged = false; if (p.packedOutput && !packedLogged) { packedLogged = true; - logf("frame %u: packed output (the output kernel stores R10G10B10A2 texels; no conversion blit)", frame); + logf("frame %u: packed output (the output kernel stores %s texels; no conversion blit)", frame, + p.packedRgba8 ? "RGBA8" : "R10G10B10A2"); } InPlaceTexture outputTexture; const size_t outputTexelBytes = p.packedOutput ? 4 : 8; @@ -9185,12 +9840,21 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand const auto depthState = static_cast(env_uint("D4R_SHIM_DEPTH_STATE", 0x40)); const auto outputState = static_cast(env_uint("D4R_SHIM_OUTPUT_STATE", 0x8)); const auto inputRecordStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; + static const bool timestampsWanted = env_uint("D4R_PROFILE_TIMESTAMPS", 1) != 0; + const bool timestamps = timing.enabled && timestampsWanted && p.vram && feature->split && + profile_timestamps_ready(*feature); + if (timestamps) + { + read_profile_timestamps(*feature, frame, timing); + profile_timestamp(*feature, list, frame, 0); + } if (p.vram) { const D3D12_RESOURCE_STATES states[4] = {inputState, depthState, inputState, inputState}; for (int index = 0; index < inputCount; ++index) transition(list, inputs[index], states[index], D3D12_RESOURCE_STATE_COPY_SOURCE); - const bool recorded = record_vram_inputs(*feature, list, slot, vramInputs, inputCount); + const bool fused = ensure_fused_color(*feature, slotIndex, color, vramInputs[0]); + const bool recorded = record_vram_inputs(*feature, list, slot, vramInputs, inputCount, fused); for (int index = 0; index < inputCount; ++index) transition(list, inputs[index], D3D12_RESOURCE_STATE_COPY_SOURCE, states[index]); if (!recorded) @@ -9253,6 +9917,20 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand list2->WriteBufferImmediate(2, markers, modes); list2->Release(); + if (feature->ngxGpuWait.load()) + { + p.gpuInputMarker = record_ngx_input_ready(*feature, list, slotIndex, frame); + if (!p.gpuInputMarker) + { + feature->ngxGpuWait.store(false); + logf("NGX feature %u: recording the GPU input marker failed; using CPU marker polling from now on", + feature->handle.Id); + } + } + + if (timestamps) + profile_timestamp(*feature, list, frame, 1); + // Present the most recent finished DLSS result. int latest; uint32_t colorFrame = 0; // the frame whose colour is presented below, 0 if none @@ -9268,7 +9946,9 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand } } const auto outputRecordStart = timing.enabled ? ProfileClock::now() : ProfileClock::time_point{}; - if (p.vram && feature->split) + if (p.vram && feature->split && p.gpuInputMarker && feature->gpuCompletion) + p.split = p.gpuCompletion = record_gpu_completion_wait(*feature, list, frame); + if (p.vram && feature->split && !p.gpuCompletion) p.split = SUCCEEDED(g_vk.split->SplitCommandListForExternalWait( list, reinterpret_cast(feature->splitSemaphore), frame)); if (!p.split) @@ -9276,6 +9956,8 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand p.outputInPlace = 0; // only a split frame waits for the result before the game reads its texture p.packedOutput = false; } + if (p.split && timestamps) + profile_timestamp(*feature, list, frame, 2); if (p.split) { // Runs in the second half of the list, once frame's result is in place. Split frames @@ -9323,6 +10005,8 @@ static NgxResult evaluate(ID3D12GraphicsCommandList* list, const NgxHandle* hand if (feature->rayReconstruction && !record_alpha_copyback(*feature, list, slot, outputState, colorFrame, p.split)) return NGX_FAIL_PLATFORM_ERROR; + if (p.split && timestamps) + profile_timestamp(*feature, list, frame, 3); if (timing.enabled) { timing.outputRecord = profile_ms(outputRecordStart, ProfileClock::now()); @@ -9361,6 +10045,10 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_EvaluateFeature_C(ID3D12GraphicsCommandList static void release_feature(Feature* feature) { + { + std::lock_guard lock(feature->ngxWaitMutex); + release_ngx_input_wait_locked(*feature, "the feature is being released", false); + } // A queued GPU-side input wait for a frame that never ran would stall the network's stream: release it // before the drain below waits for the worker (the worker also releases waits it sees stuck). if (feature->engineNet != nullptr) @@ -9415,6 +10103,12 @@ static void release_feature(Feature* feature) g.cu.eventDestroy(feature->profileStart); if (feature->profileEnd != nullptr) g.cu.eventDestroy(feature->profileEnd); + if (feature->profileInputsReleased != nullptr) + g.cu.eventDestroy(feature->profileInputsReleased); + if (feature->timestampHeap != nullptr) + feature->timestampHeap->Release(); + if (feature->timestampReadback != nullptr) + feature->timestampReadback->Release(); if (feature->outputReadyEvent != nullptr) g.cu.eventDestroy(feature->outputReadyEvent); if (feature->scratch != 0) @@ -9472,8 +10166,12 @@ D4R_EXPORT NgxResult NVSDK_NGX_D3D12_ReleaseFeature(NgxHandle* handle) { std::lock_guard lock(g_featuresMutex); for (Feature* candidate : g_features) - if (&candidate->handle == handle && !candidate->nativeVulkan && !candidate->retiring.exchange(true)) + if (&candidate->handle == handle && !candidate->nativeVulkan) { + // The API mutex serializes releases and release_feature removes the registry entry. An + // internal drain failure may already have marked this feature retiring; it still needs + // the caller's normal release to clean up, rather than becoming an unreleasable handle. + candidate->retiring.store(true); feature = candidate; break; } @@ -9493,6 +10191,7 @@ static void finish_shutdown() return; if (g_vk.ready) free_vram_pool(); + if (g_fusedColorReady) { g_fusedColorStage.destroy(); g_fusedColorReady = false; } if (g.ngxInitialized) { drain_pipeline(); diff --git a/tools/d4r_vulkan_staging.h b/tools/d4r_vulkan_staging.h index 01d3dab..a32c01b 100644 --- a/tools/d4r_vulkan_staging.h +++ b/tools/d4r_vulkan_staging.h @@ -16,6 +16,10 @@ struct D4rVkStaging { VkDescriptorSetLayout readSet = VK_NULL_HANDLE, writeSet = VK_NULL_HANDLE; VkPipelineLayout readLayout = VK_NULL_HANDLE, writeLayout = VK_NULL_HANDLE; VkPipeline readPipeline = VK_NULL_HANDLE; + // D4R_SHIM_GPU_COMPLETION: one storage buffer of words, push constants {index, value, limit} + VkDescriptorSetLayout waitSet = VK_NULL_HANDLE; + VkPipelineLayout waitLayout = VK_NULL_HANDLE; + VkPipeline waitPipeline = VK_NULL_HANDLE; struct OutputPipeline { VkFormat format; VkPipeline pipeline; } outputs[5] = {}; PFN_vkCreateSampler createSampler = nullptr; PFN_vkDestroySampler destroySampler = nullptr; @@ -81,7 +85,7 @@ struct D4rVkStaging { set.bindingCount = 2; set.pBindings = bindings; auto& setLayout = output ? writeSet : readSet; if (createSetLayout(d, &set, nullptr, &setLayout) != VK_SUCCESS) return false; - VkPushConstantRange range = {VK_SHADER_STAGE_COMPUTE_BIT, 0, output ? 8u : 12u}; + VkPushConstantRange range = {VK_SHADER_STAGE_COMPUTE_BIT, 0, output ? 8u : 16u}; VkPipelineLayoutCreateInfo layout = {VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO}; layout.setLayoutCount = 1; layout.pSetLayouts = &setLayout; layout.pushConstantRangeCount = 1; layout.pPushConstantRanges = ⦥ @@ -89,6 +93,44 @@ struct D4rVkStaging { } return pipeline(d4r_vk_read, sizeof(d4r_vk_read), readLayout, readPipeline); } + // Created on first use; false leaves GPU completion off. + bool wait_ready() { + if (waitPipeline) return true; + VkDescriptorSetLayoutBinding binding = {0, VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 1, VK_SHADER_STAGE_COMPUTE_BIT, nullptr}; + VkDescriptorSetLayoutCreateInfo set = {VK_STRUCTURE_TYPE_DESCRIPTOR_SET_LAYOUT_CREATE_INFO}; + set.bindingCount = 1; set.pBindings = &binding; + if (!waitSet && createSetLayout(device, &set, nullptr, &waitSet) != VK_SUCCESS) return false; + VkPushConstantRange range = {VK_SHADER_STAGE_COMPUTE_BIT, 0, 12}; + VkPipelineLayoutCreateInfo layout = {VK_STRUCTURE_TYPE_PIPELINE_LAYOUT_CREATE_INFO}; + layout.setLayoutCount = 1; layout.pSetLayouts = &waitSet; + layout.pushConstantRangeCount = 1; layout.pPushConstantRanges = ⦥ + if (!waitLayout && createLayout(device, &layout, nullptr, &waitLayout) != VK_SUCCESS) return false; + return pipeline(d4r_vk_wait, sizeof(d4r_vk_wait), waitLayout, waitPipeline); + } + bool wait_set(D4rVkStagingSet& result, VkBuffer buffer, VkDeviceSize bytes) { + if (!wait_ready()) return false; + VkDescriptorPoolSize size = {VK_DESCRIPTOR_TYPE_STORAGE_BUFFER, 1}; + VkDescriptorPoolCreateInfo pool = {VK_STRUCTURE_TYPE_DESCRIPTOR_POOL_CREATE_INFO}; + pool.maxSets = 1; pool.poolSizeCount = 1; pool.pPoolSizes = &size; + if (!result.pool && createPool(device, &pool, nullptr, &result.pool) != VK_SUCCESS) return false; + VkDescriptorSetAllocateInfo alloc = {VK_STRUCTURE_TYPE_DESCRIPTOR_SET_ALLOCATE_INFO}; + alloc.descriptorPool = result.pool; alloc.descriptorSetCount = 1; alloc.pSetLayouts = &waitSet; + if (!result.set && allocateSets(device, &alloc, &result.set) != VK_SUCCESS) return false; + VkDescriptorBufferInfo storage = {buffer, 0, bytes}; + VkWriteDescriptorSet write = {VK_STRUCTURE_TYPE_WRITE_DESCRIPTOR_SET}; + write.dstSet = result.set; write.dstBinding = 0; write.descriptorCount = 1; + write.descriptorType = VK_DESCRIPTOR_TYPE_STORAGE_BUFFER; write.pBufferInfo = &storage; + updateSets(device, 1, &write, 0, nullptr); + result.pipeline = waitPipeline; + return true; + } + void record_wait(VkCommandBuffer cmd, const D4rVkStagingSet& set, uint32_t index, uint32_t value, uint32_t limit) { + bindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, waitPipeline); + bindSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, waitLayout, 0, 1, &set.set, 0, nullptr); + const uint32_t constants[] = {index, value, limit}; + push(cmd, waitLayout, VK_SHADER_STAGE_COMPUTE_BIT, 0, 12, constants); + dispatch(cmd, 1, 1, 1); + } VkPipeline output_pipeline(VkFormat format) { for (auto& item : outputs) if (item.pipeline && item.format == format) return item.pipeline; const uint32_t* code = nullptr; size_t bytes = 0; @@ -135,12 +177,12 @@ struct D4rVkStaging { updateSets(device, 2, writes, 0, nullptr); return true; } - void record(VkCommandBuffer cmd, const D4rVkStagingSet& set, uint32_t width, uint32_t height, uint32_t inputMode = 0) { + void record(VkCommandBuffer cmd, const D4rVkStagingSet& set, uint32_t width, uint32_t height, uint32_t inputMode = 0, uint32_t stride = 0) { const auto layout = set.output ? writeLayout : readLayout; bindPipeline(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, set.pipeline); bindSets(cmd, VK_PIPELINE_BIND_POINT_COMPUTE, layout, 0, 1, &set.set, 0, nullptr); - const uint32_t dimensions[] = {width, height, inputMode}; - push(cmd, layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, set.output ? 8 : 12, dimensions); + const uint32_t dimensions[] = {width, height, inputMode, stride ? stride : width}; + push(cmd, layout, VK_SHADER_STAGE_COMPUTE_BIT, 0, set.output ? 8 : 16, dimensions); dispatch(cmd, (width+7)/8, (height+7)/8, 1); } void free(D4rVkStagingSet& set) { @@ -149,6 +191,9 @@ struct D4rVkStaging { } void destroy() { if (readPipeline) destroyPipeline(device, readPipeline, nullptr); + if (waitPipeline) destroyPipeline(device, waitPipeline, nullptr); + if (waitLayout) destroyLayout(device, waitLayout, nullptr); + if (waitSet) destroySetLayout(device, waitSet, nullptr); for (auto& item : outputs) if (item.pipeline) destroyPipeline(device, item.pipeline, nullptr); if (readLayout) destroyLayout(device, readLayout, nullptr); if (writeLayout) destroyLayout(device, writeLayout, nullptr); diff --git a/tools/wine_nvcuda_bridge.c b/tools/wine_nvcuda_bridge.c index c07b6e4..4681dd7 100644 --- a/tools/wine_nvcuda_bridge.c +++ b/tools/wine_nvcuda_bridge.c @@ -11,6 +11,7 @@ #include #include #include +#include #include "d4r_native_selection.h" #include "d4r_motion_dilation.h" #ifdef D4R_MOTION_KERNELS @@ -366,11 +367,41 @@ typedef struct char name[128]; uint64_t hash; int linked; + int rgba8; } NativeKernel; static NativeKernel* native_kernels; static size_t native_kernel_count; static char native_source[1024]; + +/* Packed stores are a code-object capability, not a property of the kernel's name. A legacy native + kernel ignores the R4 redirect tag. The sidecar must therefore name the exact bytes installed + alongside the PTX compatibility manifest; the per-process cache retains those same bytes. */ +static int native_rgba8_verified(const char* directory, const char* name) +{ + char path[1400], row[256], entry[128]; + unsigned long long expected = 0; + snprintf(path, sizeof(path), "%s/d4r-rgba8.txt", directory); + FILE* manifest = fopen(path, "r"); + if (manifest == NULL) return 0; + int found = 0; + while (fgets(row, sizeof(row), manifest) != NULL) + if (sscanf(row, "%127s %llx", entry, &expected) == 2 && strcmp(entry, name) == 0) + { found = 1; break; } + fclose(manifest); + if (!found) return 0; + snprintf(path, sizeof(path), "%s/%s.hsaco", directory, name); + FILE* object = fopen(path, "rb"); + if (object == NULL) return 0; + uint64_t hash = UINT64_C(14695981039346656037); + unsigned char bytes[4096]; + size_t count; + while ((count = fread(bytes, 1, sizeof(bytes), object)) != 0) + for (size_t i = 0; i < count; ++i) hash = (hash ^ bytes[i]) * UINT64_C(1099511628211); + const int valid = !ferror(object) && hash == (uint64_t)expected; + fclose(object); + return valid; +} static char native_served[1024]; static pthread_mutex_t native_lock = PTHREAD_MUTEX_INITIALIZER; @@ -506,6 +537,7 @@ static void prepare_native_kernels(const char* cache_home) snprintf(kernel->name, sizeof(kernel->name), "%s", name); kernel->hash = hash; kernel->linked = 0; + kernel->rgba8 = native_rgba8_verified(source, name); } fclose(manifest); if (!d4r_native_swin_encoders_enabled()) @@ -733,9 +765,22 @@ static struct { CUfunction function; int native; + int rgba8; } output_kernels[OUTPUT_KERNEL_CAPACITY]; /* the latest lookups, newest last */ static pthread_mutex_t output_kernel_lock = PTHREAD_MUTEX_INITIALIZER; static int last_output_native = -1; +static int last_output_rgba8; +/* Most NGX launches are network/input kernels. A monotonic membership filter lets those bypass the + output registry's lock and scan. Collisions only take the old lookup path; a bit is never cleared, + so function-handle reuse, feature recreation and concurrent lookups cannot hide an output kernel. */ +static uint64_t output_kernel_bits[16]; + +static unsigned output_kernel_bit(CUfunction function) +{ + uint64_t value = (uint64_t)(uintptr_t)function >> 4; + value ^= value >> 17; + return (unsigned)((value * UINT64_C(0x9e3779b97f4a7c15)) >> 54); +} static void note_function_lookup(CUfunction function, const char* name) { @@ -755,20 +800,34 @@ static void note_function_lookup(CUfunction function, const char* name) memmove(output_kernels, output_kernels + 1, sizeof(output_kernels) - sizeof(output_kernels[0])); output_kernels[OUTPUT_KERNEL_CAPACITY - 1].function = function; output_kernels[OUTPUT_KERNEL_CAPACITY - 1].native = native; + int rgba8 = 0; + for (size_t i = 0; i < native_kernel_count; ++i) + if (strcmp(native_kernels[i].name, name) == 0) { rgba8 = native_kernels[i].rgba8; break; } + output_kernels[OUTPUT_KERNEL_CAPACITY - 1].rgba8 = native && rgba8; pthread_mutex_unlock(&output_kernel_lock); + const unsigned bit = output_kernel_bit(function); + __atomic_fetch_or(&output_kernel_bits[bit >> 6], UINT64_C(1) << (bit & 63), __ATOMIC_RELEASE); } -static void note_launch(CUfunction function) +/* Returns 1 + native + 2 * rgba8 when `function` is an output kernel, 0 otherwise. */ +static int note_launch(CUfunction function) { + const unsigned bit = output_kernel_bit(function); + if ((__atomic_load_n(&output_kernel_bits[bit >> 6], __ATOMIC_ACQUIRE) & (UINT64_C(1) << (bit & 63))) == 0) + return 0; + int kind = 0; pthread_mutex_lock(&output_kernel_lock); /* newest first, so a recycled handle resolves with its latest lookup */ for (int i = OUTPUT_KERNEL_CAPACITY - 1; i >= 0 && output_kernels[i].function != NULL; --i) if (output_kernels[i].function == function) { __atomic_store_n(&last_output_native, output_kernels[i].native, __ATOMIC_RELAXED); + __atomic_store_n(&last_output_rgba8, output_kernels[i].rgba8, __ATOMIC_RELAXED); + kind = 1 + output_kernels[i].native + 2 * output_kernels[i].rgba8; break; } pthread_mutex_unlock(&output_kernel_lock); + return kind; } /* 1 when the last output kernel launched was d4r's native one, 0 when it was ZLUDA's compile of @@ -778,6 +837,11 @@ int WINAPI d4rOutputKernelNative(void) return __atomic_load_n(&last_output_native, __ATOMIC_RELAXED); } +int WINAPI d4rOutputKernelRgba8(void) +{ + return __atomic_load_n(&last_output_rgba8, __ATOMIC_RELAXED); +} + static void load_zluda(void) { /* Explicit experimental backend. Never silently mix module/context/array @@ -964,7 +1028,7 @@ CUresult WINAPI d4rDilateMotionVectors(const D4rMotionDilationParams* parameters D4rMotionDilationParams copy = *parameters; void* arguments[] = {©}; const uint64_t pixels = (uint64_t)copy.motion_width * copy.motion_height; - return launch(function, (unsigned)((pixels + 127) / 128), 1, 1, 128, 1, 1, 0, NULL, arguments, NULL); + return launch(function, (unsigned)((pixels + 127) / 128), 1, 1, 128, 1, 1, 0, d4r_ngx_stream(NULL), arguments, NULL); #else return CUDA_ERROR_NOT_SUPPORTED; #endif @@ -1127,6 +1191,16 @@ static void ensure_context(void) pthread_once(&context_once, create_default_zluda_context); } +/* D4R_NGX_STREAM_PRIORITY=high|normal (default off): the null-stream work of NGX, the shim and this bridge + (kernels, async copies and memsets, events, stream waits, input waits, completion callbacks) moves from + HIP's null stream to one blocking HIP stream. "high" creates it at the device's greatest priority, so its + HSA queue (and the hardware queue it maps to) dispatches ahead of normal-priority queues; "normal" creates + it at the default priority, the control that isolates the stream change itself. Every caller passes NULL + and gets the same stream, so the order of the work is unchanged; being blocking, the stream also stays + ordered with ZLUDA's own synchronous null-stream operations. Off: NULL stays HIP's null stream. */ +static CUstream d4r_ngx_stream(CUstream stream); +static void* hip_symbol(const char* name); + /* Opt-in stage instrumentation. With D4R_CUDA_LAUNCH_STATS=1 every cuLaunchKernel is followed by a context synchronize and a summary of each device allocation, surface object and texture object named in the packed @@ -1402,15 +1476,17 @@ static CUresult write_redirect_tail(CUsurfObject object, CUdeviceptr pointer, ui static unsigned int next; unsigned char* tail = ring[next++ % 256]; // dword 21: 0x5232 ('R2', RGBA16F texels) << 16 | pitch / 8 (0 = off); dwords 22-23: pointer. - // Row pitches are multiples of 8, so bit 0 of `pitch` selects 'R3': R10G10B10A2_UNORM texels. - const uint32_t word = pitch != 0 ? ((pitch & 1u) != 0 ? 0x52330000u : 0x52320000u) | ((pitch >> 3) & 0xffffu) : 0u; + // Row pitches are multiples of 8: bit 0 selects R3 (RGB10A2), bit 1 selects R4 (RGBA8). + const uint32_t word = pitch != 0 ? ((pitch & 2u) != 0 ? 0x52340000u : + (pitch & 1u) != 0 ? 0x52330000u : 0x52320000u) | + ((pitch >> 3) & 0xffffu) : 0u; memset(tail, 0, 16); memcpy(tail, &word, 4); memcpy(tail + 4, &pointer, 8); CUMEMCPYHTODASYNC_FN function = (CUMEMCPYHTODASYNC_FN)find_zluda_symbol("cuMemcpyHtoDAsync_v2"); if (function == NULL) return CUDA_ERROR_NOT_SUPPORTED; - CUresult result = function((CUdeviceptr)object + 84, tail, 12, NULL); + CUresult result = function((CUdeviceptr)object + 84, tail, 12, d4r_ngx_stream(NULL)); if (result == CUDA_SUCCESS && !async) { CUCTX_SYNCHRONIZE_FN sync = (CUCTX_SYNCHRONIZE_FN)find_zluda_symbol("cuCtxSynchronize"); @@ -2431,6 +2507,33 @@ CUresult WINAPI cuModuleGetFunction(CUfunction* function_out, CUmodule module, c return result; } +/* D4R_SHIM_GPU_COMPLETION, early form: the shim arms a completion word before NGX's evaluate. When the output + kernel that follows is d4r's native one (which writes the game-visible buffer directly; RGBA8-capable when + the frame needs the packed store), the word is queued right behind that kernel, ahead of NGX's trailing + bookkeeping kernels, which the game never reads. Any other output kernel leaves it to the shim's own write + after the output copy. Only the CUDA worker evaluates, so the arm state needs no lock. */ +static CUdeviceptr ngx_completion_pointer; +static uint32_t ngx_completion_value; +static int ngx_completion_rgba8, ngx_completion_fired; + +CUresult WINAPI d4rNgxArmCompletion(CUdeviceptr pointer, uint32_t value, int rgba8) +{ + ngx_completion_pointer = pointer; + ngx_completion_value = value; + ngx_completion_rgba8 = rgba8; + ngx_completion_fired = 0; + return CUDA_SUCCESS; +} + +/* 1 when the armed word was queued behind the native output kernel; disarms. */ +int WINAPI d4rNgxCompletionFired(void) +{ + const int fired = ngx_completion_fired; + ngx_completion_pointer = 0; + ngx_completion_fired = 0; + return fired; +} + CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x, unsigned int grid_y, unsigned int grid_z, unsigned int block_x, unsigned int block_y, unsigned int block_z, unsigned int shared_bytes, CUstream stream, @@ -2439,13 +2542,11 @@ CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x, static CULAUNCHKERNEL_FN function; if (function == NULL) function = (CULAUNCHKERNEL_FN)find_zluda_symbol("cuLaunchKernel"); - unsigned int sequence = 0; - pthread_mutex_lock(&instrumentation_lock); - sequence = ++launch_sequence; - pthread_mutex_unlock(&instrumentation_lock); + stream = d4r_ngx_stream(stream); + const unsigned int sequence = __atomic_add_fetch(&launch_sequence, 1, __ATOMIC_RELAXED); replay_dump_launch(sequence, function_handle, grid_x, grid_y, grid_z, block_x, block_y, block_z, shared_bytes, extra); - note_launch(function_handle); + const int output_kind = note_launch(function_handle); PendingKernelProfile* profile = function != NULL ? begin_kernel_profile(sequence, function_handle, stream) : NULL; CUfunction native_function = microcuda_k_lookup(function_handle); @@ -2468,6 +2569,17 @@ CUresult WINAPI cuLaunchKernel(CUfunction function_handle, unsigned int grid_x, tracef("microcuda-k prep stats runs=%llu skips=%llu", (unsigned long long)prep, (unsigned long long)skips); } end_kernel_profile(profile, result, stream); + if (result == CUDA_SUCCESS && ngx_completion_pointer != 0 && output_kind != 0) + { + const int native = ((output_kind - 1) & 1) != 0, rgba8 = (output_kind - 1) >= 2; + typedef int (*WRITE_FN)(void*, void*, uint32_t, unsigned int); + static WRITE_FN write; + if (write == NULL) + write = (WRITE_FN)hip_symbol("hipStreamWriteValue32"); + if (native && (!ngx_completion_rgba8 || rgba8) && write != NULL) + ngx_completion_fired = write(stream, (void*)(uintptr_t)ngx_completion_pointer, ngx_completion_value, 0) == 0; + ngx_completion_pointer = 0; /* one output kernel per evaluate: never fire on a later one */ + } if (result != CUDA_SUCCESS || trace_verbose()) { tracef("cuLaunchKernel[%u] %s", sequence, find_function_name(function_handle)); @@ -2569,6 +2681,7 @@ CUresult WINAPI cuMemcpyHtoDAsync(CUdeviceptr destination, const void* source, s { CUMEMCPYHTODASYNC_FN function = (CUMEMCPYHTODASYNC_FN)find_zluda_symbol("cuMemcpyHtoDAsync_v2"); if (microcuda_k_write) microcuda_k_write(destination, bytes); + stream = d4r_ngx_stream(stream); CUresult result = function != NULL ? function(destination, source, bytes, stream) : missing("cuMemcpyHtoDAsync_v2"); TRACE_CALL(result, "cuMemcpyHtoDAsync->v2 dst=0x%llx src=%p bytes=%zu stream=%p result=%d", (unsigned long long)destination, source, bytes, stream, result); @@ -2906,7 +3019,6 @@ CUresult WINAPI cuTexObjectGetResourceDesc(void* descriptor, CUtexObject object) return result; } - /* Direct pass-throughs for driver APIs whose arguments are all integers or pointers, so the Windows x64 call forwards unchanged to ZLUDA (SysV). Unversioned names resolve to the versions the CUDA 12 driver returns for them. */ @@ -2935,8 +3047,8 @@ static int d4r_forward_is_teardown(const char* target) D4R_FORWARD(cuEventCreate, cuEventCreate, 2, (D4rArg a0, D4rArg a1), (a0, a1)) D4R_FORWARD(cuEventDestroy, cuEventDestroy_v2, 1, (D4rArg a0), (a0)) D4R_FORWARD(cuEventDestroy_v2, cuEventDestroy_v2, 1, (D4rArg a0), (a0)) -D4R_FORWARD(cuEventRecord, cuEventRecord, 2, (D4rArg a0, D4rArg a1), (a0, a1)) -D4R_FORWARD(cuEventRecordWithFlags, cuEventRecordWithFlags, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) +D4R_FORWARD(cuEventRecord, cuEventRecord, 2, (D4rArg a0, D4rArg a1), (a0, (D4rArg)d4r_ngx_stream((CUstream)a1))) +D4R_FORWARD(cuEventRecordWithFlags, cuEventRecordWithFlags, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, (D4rArg)d4r_ngx_stream((CUstream)a1), a2)) CUresult WINAPI cuEventSynchronize(D4rArg a0) { if (before_teardown()) @@ -2995,9 +3107,9 @@ D4R_FORWARD(cuStreamCreate, cuStreamCreate, 2, (D4rArg a0, D4rArg a1), (a0, a1)) D4R_FORWARD(cuStreamCreateWithPriority, cuStreamCreateWithPriority, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) D4R_FORWARD(cuStreamDestroy, cuStreamDestroy_v2, 1, (D4rArg a0), (a0)) D4R_FORWARD(cuStreamDestroy_v2, cuStreamDestroy_v2, 1, (D4rArg a0), (a0)) -D4R_FORWARD(cuStreamSynchronize, cuStreamSynchronize, 1, (D4rArg a0), (a0)) -D4R_FORWARD(cuStreamQuery, cuStreamQuery, 1, (D4rArg a0), (a0)) -D4R_FORWARD(cuStreamWaitEvent, cuStreamWaitEvent, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) +D4R_FORWARD(cuStreamSynchronize, cuStreamSynchronize, 1, (D4rArg a0), ((D4rArg)d4r_ngx_stream((CUstream)a0))) +D4R_FORWARD(cuStreamQuery, cuStreamQuery, 1, (D4rArg a0), ((D4rArg)d4r_ngx_stream((CUstream)a0))) +D4R_FORWARD(cuStreamWaitEvent, cuStreamWaitEvent, 3, (D4rArg a0, D4rArg a1, D4rArg a2), ((D4rArg)d4r_ngx_stream((CUstream)a0), a1, a2)) D4R_FORWARD(cuStreamGetFlags, cuStreamGetFlags, 2, (D4rArg a0, D4rArg a1), (a0, a1)) D4R_FORWARD(cuStreamGetPriority, cuStreamGetPriority, 2, (D4rArg a0, D4rArg a1), (a0, a1)) D4R_FORWARD(cuCtxGetCurrent, cuCtxGetCurrent, 1, (D4rArg a0), (a0)) @@ -3014,14 +3126,14 @@ D4R_FORWARD(cuMemGetInfo_v2, cuMemGetInfo_v2, 2, (D4rArg a0, D4rArg a1), (a0, a1 D4R_FORWARD(cuMemcpyHtoD, cuMemcpyHtoD_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) D4R_FORWARD(cuMemcpyHtoD_v2, cuMemcpyHtoD_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) D4R_FORWARD(cuMemcpyDtoD, cuMemcpyDtoD_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) -D4R_FORWARD(cuMemcpyDtoDAsync, cuMemcpyDtoDAsync_v2, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, a3)) -D4R_FORWARD(cuMemcpyDtoHAsync, cuMemcpyDtoHAsync_v2, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, a3)) +D4R_FORWARD(cuMemcpyDtoDAsync, cuMemcpyDtoDAsync_v2, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, (D4rArg)d4r_ngx_stream((CUstream)a3))) +D4R_FORWARD(cuMemcpyDtoHAsync, cuMemcpyDtoHAsync_v2, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, (D4rArg)d4r_ngx_stream((CUstream)a3))) D4R_FORWARD(cuMemsetD8, cuMemsetD8_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) D4R_FORWARD(cuMemsetD16, cuMemsetD16_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) D4R_FORWARD(cuMemsetD32, cuMemsetD32_v2, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) -D4R_FORWARD(cuMemsetD8Async, cuMemsetD8Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, a3)) -D4R_FORWARD(cuMemsetD16Async, cuMemsetD16Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, a3)) -D4R_FORWARD(cuMemsetD32Async, cuMemsetD32Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, a3)) +D4R_FORWARD(cuMemsetD8Async, cuMemsetD8Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, (D4rArg)d4r_ngx_stream((CUstream)a3))) +D4R_FORWARD(cuMemsetD16Async, cuMemsetD16Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, (D4rArg)d4r_ngx_stream((CUstream)a3))) +D4R_FORWARD(cuMemsetD32Async, cuMemsetD32Async, 4, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3), (a0, a1, a2, (D4rArg)d4r_ngx_stream((CUstream)a3))) D4R_FORWARD(cuMemsetD2D8, cuMemsetD2D8_v2, 5, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3, D4rArg a4), (a0, a1, a2, a3, a4)) D4R_FORWARD(cuMemsetD2D32, cuMemsetD2D32_v2, 5, (D4rArg a0, D4rArg a1, D4rArg a2, D4rArg a3, D4rArg a4), (a0, a1, a2, a3, a4)) D4R_FORWARD(cuFuncSetAttribute, cuFuncSetAttribute, 3, (D4rArg a0, D4rArg a1, D4rArg a2), (a0, a1, a2)) @@ -3116,6 +3228,59 @@ static void* hip_symbol(const char* name) return hip != NULL ? dlsym(hip, name) : NULL; } +static CUstream ngx_priority_stream; +static pthread_once_t ngx_priority_once = PTHREAD_ONCE_INIT; + +static void create_ngx_priority_stream(void) +{ + const char* setting = getenv("D4R_NGX_STREAM_PRIORITY"); + const int high = setting != NULL && strcmp(setting, "high") == 0; + if (!high && (setting == NULL || strcmp(setting, "normal") != 0)) + return; + typedef int (*STREAM_PRIORITY_RANGE_FN)(int*, int*); + typedef int (*STREAM_CREATE_PRIORITY_FN)(void**, unsigned int, int); + STREAM_PRIORITY_RANGE_FN range = (STREAM_PRIORITY_RANGE_FN)hip_symbol("hipDeviceGetStreamPriorityRange"); + STREAM_CREATE_PRIORITY_FN create = (STREAM_CREATE_PRIORITY_FN)hip_symbol("hipStreamCreateWithPriority"); + int least = 0, greatest = 0; + ensure_context(); + if (context_setup_result != CUDA_SUCCESS || range == NULL || create == NULL || range(&least, &greatest) != 0) + { + tracef("D4R_NGX_STREAM_PRIORITY=%s: HIP stream priorities unavailable; NGX stays on the null stream", setting); + return; + } + void* stream = NULL; + const int priority = high ? greatest : 0; + const int result = create(&stream, 0 /* hipStreamDefault: blocking */, priority); + tracef("D4R_NGX_STREAM_PRIORITY=%s: NGX stream %p priority %d (range %d..%d) result=%d", setting, stream, + priority, least, greatest, result); + if (result == 0) + ngx_priority_stream = (CUstream)stream; +} + +static CUstream d4r_ngx_stream(CUstream stream) +{ + if (stream != NULL) + return stream; + pthread_once(&ngx_priority_once, create_ngx_priority_stream); + return ngx_priority_stream; /* NULL unless D4R_NGX_STREAM_PRIORITY created one */ +} + +CUresult WINAPI d4rNgxStreamPriority(int* priority) +{ + /* For the shim's log: 1 when NGX runs on the substituted stream, with its HIP priority. */ + typedef int (*STREAM_GET_PRIORITY_FN)(void*, int*); + CUstream stream = d4r_ngx_stream(NULL); + STREAM_GET_PRIORITY_FN get = (STREAM_GET_PRIORITY_FN)hip_symbol("hipStreamGetPriority"); + if (priority == NULL) + return CUDA_ERROR_INVALID_VALUE; + *priority = 0; + if (stream == NULL) + return CUDA_ERROR_NOT_SUPPORTED; + if (get != NULL) + get(stream, priority); + return CUDA_SUCCESS; +} + CUresult WINAPI d4rImportVulkanMemory(void* client_device, uint64_t client_memory, uint64_t bytes, CUdeviceptr* pointer, void** memory) { @@ -3282,6 +3447,146 @@ typedef int (*VK_SIGNAL_SEMAPHORE_FN)(void*, const D4rVkSemaphoreSignalInfo*); typedef int (*VK_GET_SEMAPHORE_COUNTER_FN)(void*, uint64_t, uint64_t*); typedef int (*HIP_LAUNCH_HOST_FN)(void*, void (*)(void*), void*); +/* NGX's null-stream completion, independent of Wine scheduling/event polling. CPU cancellations + and native callbacks share one timeline lock so a late callback never signals a lower value. + No HIP or Wine APIs may run from the callback. The shim queues this only for the completed-prefix + head, after every output write; resource retirement waits for its native completion condition. */ +typedef struct +{ + void* device; + uint64_t semaphore, value, pending; + VK_SIGNAL_SEMAPHORE_FN signal; + HIP_LAUNCH_HOST_FN launch; + pthread_mutex_t lock; + pthread_cond_t completed; + int busy, cancelled, detached; +} D4rNgxTimeline; + +int WINAPI d4rNgxTimelineSignal(void* handle, uint64_t value) +{ + D4rNgxTimeline* timeline = handle; + pthread_mutex_lock(&timeline->lock); + int result = 0; + if (value > timeline->value) + { + D4rVkSemaphoreSignalInfo info = {1000207005, NULL, timeline->semaphore, value}; + result = timeline->signal(timeline->device, &info); + if (result == 0) timeline->value = value; + } + pthread_mutex_unlock(&timeline->lock); + return result; +} + +static void d4r_ngx_completed(void* data) +{ + D4rNgxTimeline* timeline = data; + pthread_mutex_lock(&timeline->lock); + if (!__atomic_load_n(&timeline->cancelled, __ATOMIC_ACQUIRE) && timeline->pending > timeline->value) + { + D4rVkSemaphoreSignalInfo info = {1000207005, NULL, timeline->semaphore, timeline->pending}; + if (timeline->signal(timeline->device, &info) == 0) timeline->value = timeline->pending; + /* A failed signal leaves the worker's normal retirement to retry. */ + } + timeline->busy = 0; + pthread_cond_signal(&timeline->completed); + const int detached = timeline->detached; + pthread_mutex_unlock(&timeline->lock); + if (detached) { pthread_cond_destroy(&timeline->completed); pthread_mutex_destroy(&timeline->lock); free(timeline); } +} + +CUresult WINAPI d4rNgxTimelineCreate(void* client_device, uint64_t client_semaphore, void** handle) +{ + if (client_device == NULL || client_semaphore == 0 || handle == NULL) return CUDA_ERROR_INVALID_VALUE; + const uint64_t* device_object = (const uint64_t*)(uintptr_t)((const uint64_t*)client_device)[1]; + const uint64_t* semaphore_object = (const uint64_t*)(uintptr_t)client_semaphore; + if (device_object == NULL || device_object[1] != (uint64_t)(uintptr_t)client_device || + semaphore_object[1] != client_semaphore) return CUDA_ERROR_NOT_SUPPORTED; + void* vulkan = dlopen("libvulkan.so.1", RTLD_NOW | RTLD_NOLOAD); + VK_GET_DEVICE_PROC_ADDR_FN get = vulkan != NULL ? dlsym(vulkan, "vkGetDeviceProcAddr") : NULL; + HIP_LAUNCH_HOST_FN launch = (HIP_LAUNCH_HOST_FN)hip_symbol("hipLaunchHostFunc"); + if (get == NULL || launch == NULL) return CUDA_ERROR_NOT_SUPPORTED; + void* device = (void*)(uintptr_t)device_object[0]; + VK_SIGNAL_SEMAPHORE_FN signal = (VK_SIGNAL_SEMAPHORE_FN)get(device, "vkSignalSemaphore"); + VK_GET_SEMAPHORE_COUNTER_FN counter = (VK_GET_SEMAPHORE_COUNTER_FN)get(device, "vkGetSemaphoreCounterValue"); + uint64_t value = 0; + if (signal == NULL || counter == NULL || counter(device, semaphore_object[0], &value) != 0) + return CUDA_ERROR_NOT_SUPPORTED; + D4rNgxTimeline* timeline = calloc(1, sizeof(*timeline)); + if (timeline == NULL) return CUDA_ERROR_OUT_OF_MEMORY; + if (pthread_mutex_init(&timeline->lock, NULL) != 0) { free(timeline); return CUDA_ERROR_UNKNOWN; } + pthread_condattr_t attr; + if (pthread_condattr_init(&attr) != 0) { pthread_mutex_destroy(&timeline->lock); free(timeline); return CUDA_ERROR_UNKNOWN; } + const int condition = pthread_condattr_setclock(&attr, CLOCK_MONOTONIC) != 0 || + pthread_cond_init(&timeline->completed, &attr) != 0; + pthread_condattr_destroy(&attr); + if (condition) { pthread_mutex_destroy(&timeline->lock); free(timeline); return CUDA_ERROR_UNKNOWN; } + timeline->device = device; timeline->semaphore = semaphore_object[0]; timeline->value = value; + timeline->signal = signal; timeline->launch = launch; *handle = timeline; + return CUDA_SUCCESS; +} + +CUresult WINAPI d4rNgxTimelinePrepare(void* handle) +{ + D4rNgxTimeline* timeline = handle; + pthread_mutex_lock(&timeline->lock); + const int busy = timeline->busy; + if (!busy) __atomic_store_n(&timeline->cancelled, 0, __ATOMIC_RELEASE); + pthread_mutex_unlock(&timeline->lock); + return busy ? CUDA_ERROR_NOT_READY : CUDA_SUCCESS; +} + +void WINAPI d4rNgxTimelineCancel(void* handle) +{ + /* Called by the independent input watchdog before its coherent release store: no queue calls/locks. */ + __atomic_store_n(&((D4rNgxTimeline*)handle)->cancelled, 1, __ATOMIC_RELEASE); +} + +CUresult WINAPI d4rNgxTimelineQueue(void* handle, uint64_t value) +{ + D4rNgxTimeline* timeline = handle; + pthread_mutex_lock(&timeline->lock); + if (timeline->busy) { pthread_mutex_unlock(&timeline->lock); return CUDA_ERROR_NOT_READY; } + timeline->pending = value; timeline->busy = 1; + pthread_mutex_unlock(&timeline->lock); + const int result = timeline->launch(d4r_ngx_stream(NULL), d4r_ngx_completed, timeline); + if (result != 0) + { + pthread_mutex_lock(&timeline->lock); timeline->busy = 0; pthread_mutex_unlock(&timeline->lock); + } + return result == 0 ? CUDA_SUCCESS : CUDA_ERROR_NOT_SUPPORTED; +} + +void WINAPI d4rNgxTimelineDestroy(void* handle) +{ + D4rNgxTimeline* timeline = handle; + pthread_mutex_lock(&timeline->lock); + __atomic_store_n(&timeline->cancelled, 1, __ATOMIC_RELEASE); + const int busy = timeline->busy; + timeline->detached = busy; + pthread_mutex_unlock(&timeline->lock); + /* Normally the native condition drained the callback. If a device error prevented that drain, retain + only this self-owned ticket until its cancelled callback returns, never a pointer into Feature. */ + if (!busy) { pthread_cond_destroy(&timeline->completed); pthread_mutex_destroy(&timeline->lock); free(timeline); } +} + +CUresult WINAPI d4rNgxTimelineWait(void* handle) +{ + D4rNgxTimeline* timeline = handle; + struct timespec deadline; + if (clock_gettime(CLOCK_MONOTONIC, &deadline) != 0) return CUDA_ERROR_UNKNOWN; + deadline.tv_sec += 5; + pthread_mutex_lock(&timeline->lock); + int result = 0; + while (timeline->busy && result == 0) + result = pthread_cond_timedwait(&timeline->completed, &timeline->lock, &deadline); + const int busy = timeline->busy; + pthread_mutex_unlock(&timeline->lock); + /* The callback runs after every output write on the null stream. Its condition is sufficient to + retire those resources, without polling/recording another GPU event. A timeout keeps the ticket + owned; the caller cancels, disables this mode and falls back to its normal output synchronization. */ + return !busy ? CUDA_SUCCESS : result == ETIMEDOUT ? CUDA_ERROR_NOT_READY : CUDA_ERROR_UNKNOWN; +} + typedef struct D4rEngineNet { int k_preset; /* D4RKHIP1: the token grid comes from the output dimensions and the cycle has eight graphs */ @@ -3333,7 +3638,6 @@ typedef int (*HIP_CAPTURE_BEGIN_FN)(void*, int); typedef int (*HIP_CAPTURE_END_FN)(void*, void**); typedef int (*HIP_GRAPH_INSTANTIATE_FN)(void**, void*, void**, char*, size_t); - CUresult WINAPI d4rEngineNetReleaseInput(void* handle, uint32_t value) { D4rEngineNet* net = (D4rEngineNet*)handle; @@ -4175,7 +4479,7 @@ CUresult WINAPI d4rMemcpy2DAsync(const D4rMemcpy2D* copy, CUstream stream) D4rMemcpy2D hip = *copy; hip.srcMemoryType = (unsigned int)src; hip.dstMemoryType = (unsigned int)dst; - const int result = function(&hip, NULL); + const int result = function(&hip, d4r_ngx_stream(NULL)); if (result != 0 || trace_verbose()) tracef("d4rMemcpy2DAsync %zux%zu result=%d", copy->WidthInBytes, copy->Height, result); return result == 0 ? CUDA_SUCCESS : CUDA_ERROR_INVALID_VALUE; @@ -4194,12 +4498,68 @@ CUresult WINAPI d4rStreamWaitValue32(CUdeviceptr pointer, uint32_t value) function = (HIP_STREAM_WAIT_VALUE32_FN)hip_symbol("hipStreamWaitValue32"); if (function == NULL) return CUDA_ERROR_NOT_SUPPORTED; - const int result = function(NULL, (void*)(uintptr_t)pointer, value, 0 /* hipStreamWaitValueGte */, 0xffffffffu); + const int result = function(d4r_ngx_stream(NULL), (void*)(uintptr_t)pointer, value, 0 /* hipStreamWaitValueGte */, + 0xffffffffu); if (result != 0 || trace_verbose()) tracef("d4rStreamWaitValue32 %p >= %u result=%d", (void*)(uintptr_t)pointer, value, result); return result == 0 ? CUDA_SUCCESS : CUDA_ERROR_INVALID_VALUE; } +/* Exact per-slot tags, rather than the engine's monotonic Gte word: an old or later frame must never + satisfy NGX's wait, including across u32 wrap. ZLUDA uses HIP's null stream for null CUstream kernel + launches; d4rMemcpy2DAsync uses that same stream. This export therefore orders both input copies and + NGX kernels without adding an unsupported CUDA driver entry point to ZLUDA. The caller owns a coherent + CPU mapping and arms its watchdog before entering here, so a blocked HIP call can still be released. */ +CUresult WINAPI d4rStreamWaitValue32Eq(CUdeviceptr pointer, uint32_t value) +{ + typedef int (*GET_ATTRIBUTE_FN)(int*, int, int); + typedef int (*GET_DEVICE_FN)(int*); + ensure_context(); + if (context_setup_result != CUDA_SUCCESS) + return context_setup_result; + HIP_STREAM_WAIT_VALUE32_FN wait = (HIP_STREAM_WAIT_VALUE32_FN)hip_symbol("hipStreamWaitValue32"); + GET_ATTRIBUTE_FN attribute = (GET_ATTRIBUTE_FN)hip_symbol("hipDeviceGetAttribute"); + GET_DEVICE_FN device = (GET_DEVICE_FN)hip_symbol("hipGetDevice"); + int current = 0, supported = 0; + if (pointer == 0 || wait == NULL || attribute == NULL || device == NULL || device(¤t) != 0 || + attribute(&supported, 10013 /* hipDeviceAttributeCanUseStreamWaitValue */, current) != 0 || !supported) + return CUDA_ERROR_NOT_SUPPORTED; + const int result = wait(d4r_ngx_stream(NULL), (void*)(uintptr_t)pointer, value, 1 /* hipStreamWaitValueEq */, + 0xffffffffu); + if (result != 0 || trace_verbose()) + tracef("d4rStreamWaitValue32Eq %p == %u result=%d", (void*)(uintptr_t)pointer, value, result); + return result == 0 ? CUDA_SUCCESS : CUDA_ERROR_NOT_SUPPORTED; /* failure: no wait queued */ +} + +double WINAPI d4rThreadCpuTimeMs(void) +{ + struct timespec resolution, now; + if (clock_getres(CLOCK_THREAD_CPUTIME_ID, &resolution) != 0 || resolution.tv_sec != 0 || + resolution.tv_nsec > 100000 /* 100 us: too coarse for these sub-ms submission/wait costs */ || + clock_gettime(CLOCK_THREAD_CPUTIME_ID, &now) != 0) + return -1.0; + return (double)now.tv_sec * 1000.0 + (double)now.tv_nsec / 1000000.0; +} + +/* D4R_SHIM_GPU_COMPLETION: queues a 32-bit write of `value` to `pointer` on the stream NGX runs on, behind the + frame's output writes. The game's queue polls that word (kernels/support/vulkan_wait.comp) instead of waiting + for a CPU-signalled timeline value, so DLSS's completion reaches the game without a host round trip. */ +CUresult WINAPI d4rNgxWriteValue32(CUdeviceptr pointer, uint32_t value) +{ + static HIP_STREAM_WRITE_VALUE32_FN write; + ensure_context(); + if (context_setup_result != CUDA_SUCCESS) + return context_setup_result; + if (write == NULL) + write = (HIP_STREAM_WRITE_VALUE32_FN)hip_symbol("hipStreamWriteValue32"); + if (write == NULL || pointer == 0) + return CUDA_ERROR_NOT_SUPPORTED; + const int result = write(d4r_ngx_stream(NULL), (void*)(uintptr_t)pointer, value, 0); + if (result != 0 || trace_verbose()) + tracef("d4rNgxWriteValue32 %p = %u result=%d", (void*)(uintptr_t)pointer, value, result); + return result == 0 ? CUDA_SUCCESS : CUDA_ERROR_NOT_SUPPORTED; +} + CUresult WINAPI d4rWriteValue32(CUdeviceptr pointer, uint32_t value) { static HIP_STREAM_WRITE_VALUE32_FN write; diff --git a/tools/wine_nvcuda_bridge.spec b/tools/wine_nvcuda_bridge.spec index 01da021..573697d 100644 --- a/tools/wine_nvcuda_bridge.spec +++ b/tools/wine_nvcuda_bridge.spec @@ -104,10 +104,24 @@ @ stdcall d4rSetShutdownHook(ptr) d4rSetShutdownHook @ stdcall d4rEventSynchronize(ptr) d4rEventSynchronize @ stdcall d4rStreamWaitValue32(int64 long) d4rStreamWaitValue32 +@ stdcall d4rStreamWaitValue32Eq(int64 long) d4rStreamWaitValue32Eq +@ stdcall d4rNgxStreamPriority(ptr) d4rNgxStreamPriority +@ stdcall d4rNgxWriteValue32(int64 long) d4rNgxWriteValue32 +@ stdcall d4rNgxArmCompletion(int64 long long) d4rNgxArmCompletion +@ stdcall d4rNgxCompletionFired() d4rNgxCompletionFired +@ stdcall d4rThreadCpuTimeMs() d4rThreadCpuTimeMs @ stdcall d4rWriteValue32(int64 long) d4rWriteValue32 @ stdcall d4rSetArrayRedirect(ptr int64 long) d4rSetArrayRedirect @ stdcall d4rRegisterLinearTexture(int64 int64 int64 long long) d4rRegisterLinearTexture @ stdcall d4rSetEnv(str str long) d4rSetEnv @ stdcall d4rLoadError() d4rLoadError @ stdcall d4rOutputKernelNative() d4rOutputKernelNative +@ stdcall d4rOutputKernelRgba8() d4rOutputKernelRgba8 +@ stdcall d4rNgxTimelineCreate(ptr int64 ptr) d4rNgxTimelineCreate +@ stdcall d4rNgxTimelinePrepare(ptr) d4rNgxTimelinePrepare +@ stdcall d4rNgxTimelineCancel(ptr) d4rNgxTimelineCancel +@ stdcall d4rNgxTimelineQueue(ptr int64) d4rNgxTimelineQueue +@ stdcall d4rNgxTimelineSignal(ptr int64) d4rNgxTimelineSignal +@ stdcall d4rNgxTimelineDestroy(ptr) d4rNgxTimelineDestroy +@ stdcall d4rNgxTimelineWait(ptr) d4rNgxTimelineWait @ stdcall d4rDilateMotionVectors(ptr) d4rDilateMotionVectors