From bf1de8f40130af5d3b490ee40dcc2fb20d460a14 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 20:53:15 -0400 Subject: Switch 2: suspend channels on semaphore acquires; fix pushbuffer desyncs NVN2 gates command buffers with semaphores the CPU releases after writing them. Eden read past unmet SEMAPHORED acquires, so it executed stale commands and desynchronised. Channels now suspend on an unmet acquire and resume from the same word once it is met (scheduler keeps the waiting channels, the GPU thread polls them), as in changes.patch. Also: methods split across GPFIFO entries around a called command buffer are set aside and resumed; subchannel 2 defaults to inline-to-memory; unbound subchannels and out-of-range compute/DMA methods are dropped instead of crashing; 4-byte semaphore releases no longer write 16 bytes; Ampere SEM_PAYLOAD_HI/SEM_EXECUTE no longer spin as Kepler acquires; unmapped GPU reads log instead of asserting. Debug aids: EDEN_NX2_PB_CHECK (pushbuffer changed after submit), EDEN_NX2_QUERY_TRACE, a one-time pushbuffer dump. --- src/video_core/query_cache/query_cache.h | 5 +++++ 1 file changed, 5 insertions(+) (limited to 'src/video_core/query_cache/query_cache.h') diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h index b8c1bc4f45..11b2f631a5 100644 --- a/src/video_core/query_cache/query_cache.h +++ b/src/video_core/query_cache/query_cache.h @@ -7,6 +7,7 @@ #pragma once #include +#include #include #include #include @@ -233,6 +234,10 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type QueryPropertiesFlags flags, u32 payload, u32 subreport) { const bool has_timestamp = True(flags & QueryPropertiesFlags::HasTimeout); const bool is_fence = True(flags & QueryPropertiesFlags::IsAFence); + if (static const bool trace = std::getenv("EDEN_NX2_QUERY_TRACE") != nullptr; trace) { + LOG_INFO(HW_GPU, "nx2 report addr={:#x} type={} flags={:#x} payload={:#x}", addr, + static_cast(counter_type), static_cast(flags), payload); + } size_t streamer_id = static_cast(counter_type); auto* streamer = impl->streamers[streamer_id]; if (streamer == nullptr) [[unlikely]] { -- cgit v1.2.3 From a5e528e2d0bc784ebf24850e5913cece912c685e Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 00:24:14 -0400 Subject: nx2: get title 4 past the difficulty confirm - MME: implement MME_DMA_READ / MME_DMA_READ_FIFOED (Turing macros read query results into their load FIFO; the conditional-render macro at 0x142 needs it). - vi2 present waits up to 4 s for the frame fence (was 250 ms) so shader compile stalls don't let the game lap its equal-value gate semaphores. - Native QMD: fire the third release slot at word 52, which NVN2 patches for queue-to-queue syncs (compute queue releasing a gate the 3D queue waits on). - Compute class SET_REPORT_SEMAPHORE (0x6c0-0x6c3) releases. - Ampere lowering: AOFFI, any LOP3 truth table, STG RZ, LDG 0x981, IABS, I2FP rounding, IMAD negated addend, SHFL 0x989 (all 77 title-4 shaders lower). - Debug: EDEN_NX2_SEM_WATCH= traces acquires, releases, method arguments and compute macro emits for one semaphore block; EDEN_NX2_KEY_SCRIPT scripted key presses; per-submit logs moved to TRACE. --- src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp | 4 +- src/core/hle/service/nx2_gsv.h | 10 +- .../frontend/ampere/lower_to_maxwell.h | 144 ++++++++++++++------- src/video_core/dma_pusher.cpp | 28 ++++ src/video_core/dma_pusher.h | 1 + src/video_core/engines/kepler_compute.cpp | 63 +++++++++ src/video_core/engines/kepler_compute.h | 1 + src/video_core/engines/maxwell_3d.h | 5 + src/video_core/engines/puller.cpp | 22 +++- src/video_core/macro.cpp | 35 ++++- src/video_core/macro_turing.h | 4 +- src/video_core/native_qmd.h | 20 +-- src/video_core/query_cache/query_cache.h | 7 +- src/video_core/renderer_vulkan/vk_rasterizer.cpp | 8 ++ src/yuzu_cmd/yuzu.cpp | 46 +++++++ 15 files changed, 334 insertions(+), 64 deletions(-) (limited to 'src/video_core/query_cache/query_cache.h') diff --git a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp index cf3948c321..2a144cc7ab 100644 --- a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp +++ b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp @@ -421,7 +421,7 @@ NvResult nvhost_gpu::SubmitGPFIFOImpl(IoctlSubmitGpfifo& params, Tegra::CommandL if (g_nx2_driver) { for (const auto& entry : entries.command_lists) { - LOG_DEBUG(Service_NVDRV, "nx2 GPFIFO entry raw={:#018x} addr={:#x} size={:#x}", entry.raw, + LOG_TRACE(Service_NVDRV, "nx2 GPFIFO entry raw={:#018x} addr={:#x} size={:#x}", entry.raw, entry.addr.Value(), entry.size.Value()); } static const bool check = std::getenv("EDEN_NX2_PB_CHECK") != nullptr; @@ -471,7 +471,7 @@ NvResult nvhost_gpu::SubmitGPFIFOImpl(IoctlSubmitGpfifo& params, Tegra::CommandL } if (g_nx2_driver) { - LOG_DEBUG(Service_NVDRV, + LOG_TRACE(Service_NVDRV, "nx2 SubmitGPFIFO entries={} flags={:#x} fence_out=({}, {}) guest_value={}", params.num_entries, flags.raw, params.fence.id, params.fence.value, system.Host1x().GetSyncpointManager().GetGuestSyncpointValue(params.fence.id)); diff --git a/src/core/hle/service/nx2_gsv.h b/src/core/hle/service/nx2_gsv.h index 8073e0bd96..89f3a9fd50 100644 --- a/src/core/hle/service/nx2_gsv.h +++ b/src/core/hle/service/nx2_gsv.h @@ -580,17 +580,25 @@ private: // free, so pace the game here instead: wait (bounded, in case the GPU waits on the CPU) until // Eden's GPU thread has processed the frame. Otherwise the game laps the GPU thread and // rewrites command memory it hasn't read yet. + u32 m_wait_timeouts{}; + void WaitForFrame(u32 fence_id, u32 fence_value) { if (fence_id >= 192) { return; } auto& syncpoints = system.Host1x().GetSyncpointManager(); const auto start = std::chrono::steady_clock::now(); - const auto deadline = start + std::chrono::milliseconds{250}; + // Long enough to cover shader/pipeline compile stalls: a game that gets more than a few + // command buffers ahead misses the equal-value gates its other channels acquire. + const auto deadline = start + std::chrono::milliseconds{4000}; while (!syncpoints.IsReadyGuest(fence_id, fence_value) && std::chrono::steady_clock::now() < deadline && system.IsPoweredOn()) { std::this_thread::sleep_for(std::chrono::microseconds{200}); } + if (!syncpoints.IsReadyGuest(fence_id, fence_value) && m_wait_timeouts++ < 64) { + LOG_WARNING(Service, "nx2:vi2 present #{} fence ({}, {}) not reached in 4 s (now {})", + m_presents, fence_id, fence_value, syncpoints.GetGuestSyncpointValue(fence_id)); + } if (m_presents <= 8 || (m_presents % 600) == 0) { LOG_INFO(Service, "nx2:vi2 present #{} fence ({}, {}) now {}, waited {} us", m_presents, fence_id, fence_value, syncpoints.GetGuestSyncpointValue(fence_id), diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 2c3ef28751..01dc1d8425 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -53,7 +53,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { for (const auto& instruction : code) { const u32 op = instruction.Bits(0, 12); uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0xc0c || op == 0x986 || op == 0x9a8 || op == 0xc12 || op == 0xc17 || op == 0xe24 || op == 0xc24 || op == 0xd00 || op == 0xd09 || op == 0x9c3 || op == 0xc45 || op == 0x2a4 || op == 0x891 || op == 0x890 || op == 0x290 || op == 0x28c || op == 0x88c || op == 0x287 || op == 0x887 || op == 0x958 || op == 0xab9 || op == 0xabb || op == 0x899 || op == 0xc20 || op == 0xe21 || op == 0xe23 || op == 0xc02 || - op == 0x886 || op == 0x3c2 || op == 0xc82 || op == 0x882 || op == 0x292 || op == 0x892 || op == 0x291 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || (op == 0x81c && instruction.Bits(67, 1)); + op == 0x886 || op == 0x3c2 || op == 0xc82 || op == 0x882 || op == 0x292 || op == 0x892 || op == 0x291 || op == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || (op == 0x81c && instruction.Bits(67, 1)); } const auto address_of = [](u32 index) { return 0x50U + index * 8 + ((index + 1) / 3) * 8; }; const auto fail = [&](u32 index, std::string_view reason) { @@ -454,16 +454,23 @@ inline LoweredProgram LowerToMaxwell(std::span code) { } break; } - case 0x824: { // Captured low-word IMAD forms - if (insn.words[1] == 1 && (insn.words[2] & 0xffffff00U) == 0x078e0200) { + case 0x824: { // Captured low-word IMAD forms; bit 75 negates the addend. + const bool negate_c = insn.Bits(75, 1) != 0; + if (insn.words[1] == 1 && (insn.words[2] & 0xfffff700U) == 0x078e0200) { emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | - (u64{insn.Bits(64, 8)} << 20)); + (u64{insn.Bits(64, 8)} << 20) | (u64{negate_c} << 48)); break; } if (insn.words[1] != 48 || insn.words[2] != 0x078e02ff) { - if ((insn.words[2] & ~0x2ffU) != 0x078e0000) + if ((insn.words[2] & ~0xaffU) != 0x078e0000) return fail(index, "native integer multiply mode"); - multiply_low(insn.Bits(24, 8), insn.words[1], insn.Bits(64, 8), true); + u32 c = insn.Bits(64, 8); + if (negate_c && c != 255) { + emit(0x5c10000000000000ULL | predicate | 252 | (u64{255} << 8) | (u64{c} << 20) | + (u64{1} << 48)); + c = 252; + } + multiply_low(insn.Bits(24, 8), insn.words[1], c, true); break; } const u64 source = insn.Bits(24, 8); @@ -631,6 +638,14 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{offset} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51)); break; } + case 0x989: { // SHFL, immediate lane and register clamp/segmentation mask, any mode. + if ((insn.words[1] & 0xf01fffffU) || (insn.words[2] & ~0x000e00ffU)) + return fail(index, "native register-mask shuffle mode"); + emit(base(0xef10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(53, 5)} << 20) | (u64{1} << 28) | (u64{insn.Bits(58, 2)} << 30) | + (u64{insn.Bits(64, 8)} << 39) | (u64{insn.Bits(81, 3)} << 48)); + break; + } case 0x589: { // SHFL.IDX, register lane and immediate segmentation mask. if ((insn.words[2] & ~0x000e0000U) || insn.Bits(53, 11)) return fail(index, "native shuffle mode"); @@ -882,22 +897,50 @@ inline LoweredProgram LowerToMaxwell(std::span code) { emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); break; } - case 0x381: { // LDG.E of 32/64/128 bits through a 64-bit register address. + case 0x381: case 0x981: { // LDG.E of 32/64/128 bits, register or uniform 64-bit address. const u32 size = insn.Bits(73, 3); - if ((insn.words[2] & ~0xe00U) != 0x001e8100 || size < 4 || size > 6 || - insn.Bits(32, 8) || (insn.Bits(24, 8) & 1) || insn.Bits(24, 8) >= 246 || + u32 address = insn.Bits(24, 8); + if (size < 4 || size > 6 || (opcode == 0x381 && insn.Bits(32, 8)) || (size == 5 && (dest & 1)) || (size == 6 && (dest & 3))) return fail(index, "native global load mode"); - emit(base(0xeed0000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + if (opcode == 0x381) { + if ((insn.words[2] & ~0xe00U) != 0x001e8100 || (address & 1) || address >= 246) + return fail(index, "native global load mode"); + } else { + // The uniform base sits where a register source would (bits 32-37). + const u32 uniform_address = + insn.Bits(38, 2) ? ~0U : uniform_register(insn.Bits(32, 6)); + if ((insn.words[2] & ~(0x04000000U | 0x3fU | 0xe00U | 0x2000U)) != 0x081e8100 || + uniform_address == ~0U || (uniform_address & 1)) + return fail(index, "native global load uniform addressing mode"); + if (address == 255) { + address = uniform_address; + } else { + if (insn.Bits(90, 1) && ((address & 1) || address >= 246)) + return fail(index, "native global load register address pair"); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{uniform_address} << 8) | + (u64{address} << 20) | (u64{1} << 47)); + emit(0x5c10000000000000ULL | predicate | 251 | (u64{uniform_address + 1} << 8) | + (u64{insn.Bits(90, 1) ? address + 1 : 255} << 20) | (u64{1} << 43)); + address = 250; + } + } + emit(base(0xeed0000000000000ULL) | (u64{address} << 8) | (u64{insn.Bits(40, 24)} << 20) | (u64{1} << 45) | (u64{size} << 48)); break; } case 0x386: case 0x986: { // Global store of 32/64/128 bits, 64-bit address. - const u32 data = insn.Bits(32, 8); + u32 data = insn.Bits(32, 8); const u32 size = insn.Bits(73, 3); - if (size < 4 || size > 6 || (size == 5 && (data & 1)) || (size == 6 && (data & 3)) || - data + (1U << (size - 4)) > 246) + if (size < 4 || size > 6 || + (data != 255 && ((size == 5 && (data & 1)) || (size == 6 && (data & 3)) || + data + (1U << (size - 4)) > 246))) return fail(index, "native global store size"); + if (data == 255 && size > 4) { + // Zeros wider than one register: store from zeroed scratch registers. + for (u32 i = 0; i < (1U << (size - 4)); ++i) mov_imm(250 + i, 0); + data = 250; + } u32 address = insn.Bits(24, 8); bool extended = insn.Bits(72, 1) != 0; if (opcode == 0x986) { @@ -1060,13 +1103,14 @@ inline LoweredProgram LowerToMaxwell(std::span code) { const auto general_tex = [&]() -> bool { constexpr std::array coordinate_count{1, 2, 3, 3, 2, 3, 0, 4}; constexpr std::array maxwell_type{0, 2, 4, 6, 1, 3, 0, 7}; - if (dimension == 6 || insn.Bits(76, 2) || insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || + const bool offsets = insn.Bits(76, 1) != 0; + if (dimension == 6 || insn.Bits(77, 1) || insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) >= 18 || (lod > 3 && lod != 6 && lod != 7) || (depth_compare && mask != 1) || mask == 0) return false; const u32 coordinates = coordinate_count[dimension]; const bool has_lod = lod == 2 || lod == 3 || lod == 6 || lod == 7; - const u32 count = coordinates + (has_lod ? 1 : 0) + (depth_compare ? 1 : 0); + const u32 count = coordinates + (has_lod ? 1 : 0) + (offsets ? 1 : 0) + (depth_compare ? 1 : 0); const u32 in_a = count <= 4 ? (count + 1) / 2 : 4; const auto value = [&](u32 i) -> u32 { if (i < in_a) return source_a == 255 ? 255U : source_a + i; @@ -1082,16 +1126,18 @@ inline LoweredProgram LowerToMaxwell(std::span code) { const auto copy = [&](u32 target, u32 source) { emit(0x5c98000000000000ULL | predicate | target | (u64{15} << 39) | (u64{source} << 20)); }; - // Coordinates to R250.., LOD/bias then reference to R246/R247. + // Coordinates to R250.., then LOD/bias, offsets and reference to R246... for (u32 i = 0; i < coordinates; ++i) copy(250 + i, value(i)); u32 meta = 246; - if (has_lod) copy(meta++, value(coordinates)); - if (depth_compare) copy(meta++, value(coordinates + (has_lod ? 1 : 0))); + u32 next = coordinates; + if (has_lod) copy(meta++, value(next++)); + if (offsets) copy(meta++, value(next++)); + if (depth_compare) copy(meta++, value(next++)); emit(0xc000000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{meta == 246 ? 255U : 246U} << 20) | (u64{maxwell_type[dimension]} << 28) | (u64{mask} << 31) | (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{depth_compare} << 50) | (u64{7} << 51) | - (u64{lod} << 55)); + (u64{offsets} << 54) | (u64{lod} << 55)); const u32 components = static_cast(std::popcount(mask)); if (insn.Bits(79, 1)) { // FP16 results: (c0, c1) packed into the first destination, (c2, c3) into @@ -1320,9 +1366,11 @@ inline LoweredProgram LowerToMaxwell(std::span code) { break; } case 0x245: case 0xa45: case 0xc45: { // I2FP.F32.U32 - if ((insn.words[2] & ~(0x400U | (opcode == 0xc45 ? 0x08000000U : 0U))) != 0x00201000) return fail(index, "native packed integer-to-float format"); + // Bits 78-79: rounding (RN/RM/RP/RZ), as Maxwell I2F's bits 39-40. + if ((insn.words[2] & ~(0x400U | 0xc000U | (opcode == 0xc45 ? 0x08000000U : 0U))) != 0x00201000) return fail(index, "native packed integer-to-float format"); const u32 src = operand(1, false); - emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | (u64{insn.Bits(74, 1)} << 13) | (u64{src} << 20)); + emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | (u64{insn.Bits(74, 1)} << 13) | (u64{src} << 20) | + (u64{insn.Bits(78, 2)} << 39)); break; } case 0x207: case 0x807: case 0xa07: @@ -1357,38 +1405,38 @@ inline LoweredProgram LowerToMaxwell(std::span code) { emit(out); break; } - case 0x212: case 0x812: case 0xc12: { // Bounded LOP3 bitwise forms, no predicate output + case 0x212: case 0x812: case 0xa12: case 0xc12: { // LOP3.LUT, any table, optional P = (result != 0) const u32 lut = insn.Bits(72, 8); - if ((lut != 0xc0 && lut != 0xfc && lut != 0xf8 && lut != 0x48 && lut != 0xe2) || insn.Bits(80, 1) || - insn.Bits(87, 4) != 15 || - (opcode == 0x212 && insn.Bits(40, 24))) return fail(index, "native LUT/predicate mode"); - u32 b = opcode == 0xc12 ? uniform_register(insn.Bits(32, 8)) : insn.Bits(32, 8); - if (b == ~0U) return fail(index, "native uniform logic source"); - if (opcode == 0x812) { + const u32 pred_target = insn.Bits(81, 3); + if (insn.Bits(80, 1) || insn.Bits(87, 4) != 15 || + (opcode == 0x212 && insn.Bits(40, 24)) || (opcode == 0xa12 && insn.Bits(62, 2))) + return fail(index, "native LUT/predicate mode"); + u32 b = 255; + if (opcode == 0x212) { + b = insn.Bits(32, 8); + } else if (opcode == 0x812) { mov_imm(251, insn.words[1]); b = 251; - } - const u32 a = insn.Bits(24, 8), c = insn.Bits(64, 8); - const auto logical = [&](u32 target, u32 lhs, u32 rhs, u32 operation, u32 pred_target = 7) { - emit(0x5c40000000000000ULL | predicate | target | (u64{lhs} << 8) | - (u64{rhs} << 20) | (u64{operation} << 41) | - (u64{pred_target == 7 ? 0U : 3U} << 44) | (u64{pred_target} << 48)); - }; - const u32 pred_target = insn.Bits(81, 3); - if (lut == 0xe2) { - logical(250, a, b, 0); - logical(252, c, b, 0); - result.instructions.back() |= u64{1} << 40; - logical(static_cast(dest), 250, 252, 1, pred_target); - } else if (lut == 0xf8) { - logical(250, b, c, false); - logical(static_cast(dest), a, 250, 1, pred_target); - } else if (lut == 0x48) { - logical(250, a, c, 2); - logical(static_cast(dest), 250, b, 0, pred_target); + } else if (opcode == 0xc12) { + b = uniform_register(insn.Bits(32, 8)); + if (b == ~0U) return fail(index, "native uniform logic source"); } else { - logical(static_cast(dest), a, b, lut == 0xfc ? 1U : 0U, pred_target); + b = operand(1, false); } + const u32 a = insn.Bits(24, 8), c = insn.Bits(64, 8); + emit(0x5be0000000000000ULL | predicate | dest | (u64{a} << 8) | (u64{b} << 20) | + (u64{lut} << 28) | (u64{pred_target == 7 ? 0U : 3U} << 36) | (u64{c} << 39) | + (u64{pred_target} << 48)); + break; + } + case 0x213: case 0xa13: case 0xc13: { // IABS: max(x, 0 - x), signed. + if (insn.words[2] & ~(opcode == 0xc13 ? 0x08000000U : 0U)) + return fail(index, "native integer absolute mode"); + const u32 source = operand(1, false); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{255} << 8) | (u64{source} << 20) | + (u64{1} << 48)); + emit(base(0x5c20000000000000ULL) | (u64{source} << 8) | (u64{250} << 20) | + (u64{15} << 39) | (u64{1} << 48)); break; } case 0x217: case 0x817: case 0xa17: case 0xc17: { // IMNMX register, immediate, cbuf or uniform operand. diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp index 2c95b6de03..040bf18362 100644 --- a/src/video_core/dma_pusher.cpp +++ b/src/video_core/dma_pusher.cpp @@ -13,6 +13,7 @@ #include "core/core.h" #include "core/memory.h" #include "video_core/dma_pusher.h" +#include "video_core/control/channel_state.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/gpu.h" #include "video_core/guest_memory.h" @@ -134,6 +135,16 @@ bool DmaPusher::Step() { if (++dma_pushbuffer_subindex >= command_list_size) { // A split method resumes within the submission that split it; ring addresses repeat // every frame, so a stale one could otherwise swallow the headers of a later entry. + for (const auto& parked : parked_methods) { + static std::atomic reported{}; + if (reported.fetch_add(1) < 16) { + LOG_WARNING(HW_GPU, + "Submission ended with parked method {:#x} (sub {}, {} words left, " + "resume {:#x})", + parked.method, parked.subchannel, parked.method_count, + parked.resume_address); + } + } parked_methods.clear(); dma_pushbuffer.pop(); dma_pushbuffer_subindex = 0; @@ -264,7 +275,21 @@ void DmaPusher::SetState(const CommandHeader& command_header) { dma_state.method_count = command_header.method_count; } +// EDEN_NX2_SEM_WATCH=: log every method whose argument is the low word of an +// address in the 16 bytes there (finds the engine and method that release a semaphore). +void DmaPusher::WatchArgument(u32 method, u32 argument) const { + static const u32 watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? static_cast(std::strtoull(value, nullptr, 16)) : 0U; + }(); + if (watch != 0 && (argument & ~0xfU) == watch) { + LOG_INFO(HW_GPU, "Semaphore watch: channel {} subchannel {} method {:#x} argument {:#x}", + channel_state.bind_id, dma_state.subchannel, method, argument); + } +} + void DmaPusher::CallMethod(u32 argument) { + WatchArgument(dma_state.method, argument); if (dma_state.method < non_puller_methods) { puller.CallPullerMethod(*this, Engines::Puller::MethodCall{ dma_state.method, @@ -289,6 +314,9 @@ void DmaPusher::CallMethod(u32 argument) { } void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) { + for (u32 i = 0; i < num_methods; ++i) { + WatchArgument(dma_state.method + (dma_state.non_incrementing ? 0 : i), base_start[i]); + } if (dma_state.method < non_puller_methods) { puller.CallMultiMethod(*this, dma_state.method, dma_state.subchannel, base_start, num_methods, dma_state.method_count); } else { diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h index 24fedba21b..bbcc345c84 100644 --- a/src/video_core/dma_pusher.h +++ b/src/video_core/dma_pusher.h @@ -173,6 +173,7 @@ private: void CallMethod(u32 argument); void CallMultiMethod(const u32* base_start, u32 num_methods); void ReportUnboundSubchannel(); + void WatchArgument(u32 method, u32 argument) const; public: /// Logs the GPFIFO entry being processed (and the one before it), read through the GPU MMU diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 7710ea7dd2..689f85a1d3 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include #include #include @@ -46,6 +47,46 @@ void KeplerCompute::ConsumeSinkImpl(Core::System& system) { method_sink.clear(); } +// SET_REPORT_SEMAPHORE_A-D (0x6c0-0x6c3): address high/low, payload, then operation in bits 1:0 +// (release / acquire / report only / trap), reduction enable bit 3 with the op in bits 11:9, and +// a one-word (no timestamp) structure in bit 28. NVN2 releases queue syncs from compute this way. +void KeplerCompute::ProcessNativeReportSemaphore(u32 argument) { + const GPUVAddr address = (GPUVAddr{native_registers[0x6c0] & 0x1ffffff} << 32) | + native_registers[0x6c1]; + const u32 payload = native_registers[0x6c2]; + const u32 operation = argument & 3; + const bool one_word = (argument >> 28) & 1; + static u32 logged{}; + if (logged++ < 32) { + LOG_INFO(HW_GPU, "Native compute report semaphore {:#x} = {:#x}, control {:#x}", address, + payload, argument); + } + if (operation != 0 && operation != 2) { + LOG_WARNING(HW_GPU, "Native compute report semaphore operation {} not supported", operation); + return; + } + if ((argument >> 3) & 1) { + const u32 reduction = (argument >> 9) & 7; + rasterizer->SignalFence([this, address, payload, reduction] { + const u32 old_value = memory_manager.Read(address); + const u32 value = reduction == 0 ? old_value + payload + : reduction == 1 ? std::min(old_value, payload) + : reduction == 2 ? std::max(old_value, payload) + : reduction == 3 ? (old_value >= payload ? 0 : old_value + 1) + : reduction == 4 ? (old_value == 0 || old_value > payload ? payload : old_value - 1) + : reduction == 5 ? (old_value & payload) + : reduction == 6 ? (old_value | payload) + : (old_value ^ payload); + memory_manager.Write(address, value); + }); + return; + } + rasterizer->Query(address, VideoCommon::QueryType::Payload, + one_word ? VideoCommon::QueryPropertiesFlags::IsAFence + : VideoCommon::QueryPropertiesFlags::HasTimeout, + payload, 0); +} + bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call) { if (native_context) { if (method == 0x45 || method == 0x47) native_compute_uploaded = true; @@ -65,11 +106,22 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar if (!is_last_call) return true; const u32 entry = ((*native_executing_macro - 0xe00) / 2) % native_macro_positions.size(); native_executing_macro.reset(); + static const u32 watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? static_cast(std::strtoull(value, nullptr, 16)) : 0U; + }(); + const bool watched = watch != 0 && std::ranges::any_of(native_parameters, + [](u32 parameter) { return (parameter & ~0xfU) == watch; }); + if (watched) { + LOG_INFO(HW_GPU, "Semaphore watch: compute macro start={:#x} parameters={:08x}", + native_macro_positions[entry], fmt::join(native_parameters, ",")); + } const auto result = native_macro->native_interpreter.Execute(native_macro->native_code, native_macro_positions[entry], native_parameters, [&](u32 address) { return address < native_registers.size() ? native_registers[address] : 0U; }, [&](u32 address, u32 value) { static u32 diagnostic_writes{}; + if (watched) LOG_INFO(HW_GPU, "Semaphore watch: compute macro emit {:#x} = {:#x}", address, value); if (diagnostic_writes++ < 4096) LOG_DEBUG(HW_GPU, "Native compute MME emit method={:#x}, value={:#x}", address, value); if (address >= 0xe00) return false; CallMethod(system, address, value, true); @@ -89,6 +141,10 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar return true; } native_registers[method] = argument; + if (method == 0x6c3) { // SET_REPORT_SEMAPHORE_D: the compute class's semaphore release + ProcessNativeReportSemaphore(argument); + return true; + } if (method == 0x5a || (method >= 0xc6 && method <= 0x108)) { VideoCommon::TraceNativeQuery(12, memory_manager.GetID(), method, 0, 0, argument, 0); } @@ -157,6 +213,13 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar [&](u64 address, NativeQmdScheduler::Descriptor& descriptor) { if (!memory_manager.IsFullyMappedRange(address, sizeof(descriptor))) return false; memory_manager.ReadBlockUnsafe(address, descriptor.data(), sizeof(descriptor)); + if (static u32 logged{}; (descriptor[53] & (1U << 23)) && logged < 8) { + ++logged; + LOG_INFO(HW_GPU, "Native QMD {:#x} releases: {:08x} | {:08x} | {:08x}", address, + fmt::join(std::span{descriptor}.subspan(24, 4), ","), + fmt::join(std::span{descriptor}.subspan(28, 4), ","), + fmt::join(std::span{descriptor}.subspan(52, 4), ",")); + } for (u32 slot = 0; slot < 2; ++slot) { const u32 flags = descriptor[25 + slot * 4]; if (!(flags & (1U << 23))) continue; diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index 99cf74d1e8..51d31b9c95 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -233,6 +233,7 @@ public: private: bool DispatchNativeCompute(const NativeQmdScheduler::Descriptor& descriptor); + void ProcessNativeReportSemaphore(u32 argument); bool ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call); NativeQmdScheduler native_qmd; bool native_mode{}; diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h index 8cbc4f4274..29db8022e5 100644 --- a/src/video_core/engines/maxwell_3d.h +++ b/src/video_core/engines/maxwell_3d.h @@ -3135,6 +3135,11 @@ public: /// Reads a register value located at the input method address u32 GetRegisterValue(u32 method) const; + /// GPU memory, for Turing macros that read memory into their FIFO or data RAM. + MemoryManager& GetMemoryManager() { + return memory_manager; + } + void SetNativeTuringMacroMode(bool enabled) { GetNativeMacroEngine().native_turing = enabled; upload_state.SetNativePartialTransfers(enabled); diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 97a58903fe..047e383b85 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -68,9 +68,16 @@ bool Puller::CanResumeAcquire(DmaPusher& dma_pusher) { if (!pending_acquire->reported && std::chrono::steady_clock::now() - pending_acquire->since > std::chrono::seconds{2}) { pending_acquire->reported = true; - LOG_WARNING(HW_GPU, "Channel {} waits on semaphore {:#x} (op {}, {:#x}, now {:#x}) for 2 s", + const auto device = dma_pusher.memory_manager.GpuToCpuAddress(pending_acquire->address); + LOG_WARNING(HW_GPU, "Channel {} waits on semaphore {:#x} (op {}, {:#x}, now {:#x}) for 2 s; " + "device {:#x}, buffer cache holds GPU data: {}, words {:08x} {:08x} {:08x} {:08x}", dma_pusher.channel_state.bind_id, pending_acquire->address, - pending_acquire->operation, pending_acquire->value, current); + pending_acquire->operation, pending_acquire->value, current, device.value_or(0), + device && dma_pusher.rasterizer->MustFlushRegion(*device, 16, VideoCommon::CacheType::All), + dma_pusher.memory_manager.Read(pending_acquire->address & ~GPUVAddr{0xf}), + dma_pusher.memory_manager.Read((pending_acquire->address & ~GPUVAddr{0xf}) + 4), + dma_pusher.memory_manager.Read((pending_acquire->address & ~GPUVAddr{0xf}) + 8), + dma_pusher.memory_manager.Read((pending_acquire->address & ~GPUVAddr{0xf}) + 12)); } return false; } @@ -157,6 +164,17 @@ void Puller::ProcessFenceActionMethod(DmaPusher& dma_pusher) { void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { const auto semaphoreOperationMask = 0xF; const auto op = GpuSemaphoreOperation(regs.semaphore_trigger & semaphoreOperationMask); + // EDEN_NX2_SEM_WATCH=: log every release/acquire on the 16 bytes there. + static const GPUVAddr watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? GPUVAddr{std::strtoull(value, nullptr, 16)} : GPUVAddr{}; + }(); + if (watch != 0 && (regs.semaphore_address.SemaphoreAddress() & ~GPUVAddr{0xf}) == watch) { + LOG_INFO(HW_GPU, "Semaphore watch: channel {} op {} {:#x} value {:#x} (now {:#x})", + dma_pusher.channel_state.bind_id, static_cast(op), + regs.semaphore_address.SemaphoreAddress(), regs.semaphore_sequence, + dma_pusher.memory_manager.Read(regs.semaphore_address.SemaphoreAddress())); + } if (op == GpuSemaphoreOperation::WriteLong) { const GPUVAddr sequence_address{regs.semaphore_address.SemaphoreAddress()}; const u32 payload = regs.semaphore_sequence; diff --git a/src/video_core/macro.cpp b/src/video_core/macro.cpp index a8acd81682..06c649260d 100644 --- a/src/video_core/macro.cpp +++ b/src/video_core/macro.cpp @@ -27,6 +27,7 @@ #include "video_core/engines/maxwell_3d.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/dirty_flags.h" +#include "video_core/memory_manager.h" #include "video_core/rasterizer_interface.h" #include "video_core/macro.h" @@ -1351,11 +1352,43 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u } LOG_DEBUG(HW_GPU, "Native MME context={:#x}", context); LOG_DEBUG(HW_GPU, "Native MME start={:#x}, parameters={}, data={:08x}", method, parameters.size(), fmt::join(parameters, ",")); - const auto result = native_interpreter.Execute(native_code, method, parameters, + std::vector fifo(parameters.begin(), parameters.end()); + const auto result = native_interpreter.Execute(native_code, method, fifo, [&](u32 address) { return address < Maxwell3D::Regs::NUM_REGS ? maxwell3d.GetRegisterValue(address) : 0U; }, [&](u32 address, u32 value) { if (address >= 0xe00) return false; + if (address == 0x556) { + static u32 enable_logs{}; + if (enable_logs++ < 400) LOG_INFO(HW_GPU, "Native MME start={:#x} render enable mode={}", method, value); + } maxwell3d.CallMethod(system, address, value, true); + // Turing MME memory reads (methods 0x550 SET_MME_MEM_ADDRESS_A/B, 0x558 data RAM + // address, 0x55c MME_DMA_READ into data RAM, 0x560 MME_DMA_READ_FIFOED into + // the load FIFO, which the macro then reads after a load barrier). + constexpr u32 MemAddressA = 0x154, MemAddressB = 0x155, DataRamAddress = 0x156; + constexpr u32 DmaRead = 0x157, DmaReadFifoed = 0x158; + if (address == DmaRead || address == DmaReadFifoed) { + const GPUVAddr source = (GPUVAddr{maxwell3d.GetRegisterValue(MemAddressA) & 0x1ffffff} << 32) | + maxwell3d.GetRegisterValue(MemAddressB); + auto& memory = maxwell3d.GetMemoryManager(); + const u32 ram = maxwell3d.GetRegisterValue(DataRamAddress); + static u32 dma_logs{}; + if (dma_logs < 400) { + ++dma_logs; + std::array words{}; + for (u32 i = 0; i < std::min(value, 8); ++i) words[i] = memory.Read(source + u64{i} * 4); + LOG_INFO(HW_GPU, "Native MME DMA read{} start={:#x} source={:#x} count={} words={:08x}", + address == DmaReadFifoed ? " fifoed" : "", method, source, value, fmt::join(words, ",")); + } + for (u32 i = 0; i < std::min(value, 0x1000); ++i) { + const u32 word = memory.Read(source + u64{i} * 4); + if (address == DmaReadFifoed) { + fifo.push_back(word); + } else if (ram + i < native_interpreter.dram.size()) { + native_interpreter.dram[ram + i] = word; + } + } + } return true; }); if (!result.success) { diff --git a/src/video_core/macro_turing.h b/src/video_core/macro_turing.h index fe82ce4a4f..a0a40000c2 100644 --- a/src/video_core/macro_turing.h +++ b/src/video_core/macro_turing.h @@ -10,6 +10,7 @@ #include #include #include +#include namespace Tegra { @@ -25,8 +26,9 @@ public: }; template + // `parameters` is the load FIFO; `emit` may append to it (MME_DMA_READ_FIFOED). Result Execute(std::span code, Word start, - std::span parameters, Read read, Emit emit) { + std::vector& parameters, Read read, Emit emit) { std::array regs{}; Word initialized{}, carry{}, method{}, increment{}; Word pc = start, loop_count{}, loop_start{}, loop_end{}; diff --git a/src/video_core/native_qmd.h b/src/video_core/native_qmd.h index c9ae3b8cc0..c06f14e6b2 100644 --- a/src/video_core/native_qmd.h +++ b/src/video_core/native_qmd.h @@ -61,9 +61,9 @@ public: if (!read(request.address, state.descriptor)) return "unmapped QMD"; const auto qmd = state.descriptor; if ((qmd[18] & 0xff) != 0x30 || (qmd[4] & 0x100)) return "unsupported QMD version or queue"; - // Validate both release slots before making either visible. - for (Word slot = 0; slot < 2; ++slot) { - const Word flags = qmd[25 + slot * 4]; + // Validate the release slots before making any visible. + for (const Word base : ReleaseSlots) { + const Word flags = qmd[base + 1]; if (!(flags & (1U << 23))) continue; const bool add = (flags & (1U << 26)) != 0; if ((flags & (0x3bU << 24)) || (flags >> 30) != 1 || @@ -76,15 +76,15 @@ public: if (qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff) && !dispatch(qmd)) return "native shader dispatch required"; state.completed = true; - for (Word slot = 0; slot < 2; ++slot) { - const Word flags = qmd[25 + slot * 4]; + for (const Word base : ReleaseSlots) { + const Word flags = qmd[base + 1]; if (!(flags & (1U << 23))) continue; - const Address destination = (Address{flags & 0xff} << 32) | qmd[24 + slot * 4]; + const Address destination = (Address{flags & 0xff} << 32) | qmd[base]; bool released{}; if constexpr (std::is_invocable_r_v) { - released = release(destination, qmd[26 + slot * 4], flags); + released = release(destination, qmd[base + 2], flags); } else { - released = release(destination, qmd[26 + slot * 4]); + released = release(destination, qmd[base + 2]); } if (!released) return "unmapped QMD release"; } @@ -99,6 +99,10 @@ public: return {}; } private: + // Release slots (address low, flags | address high, payload). NVN2 patches a third one at + // word 52 for queue syncs (e.g. a compute queue releasing a semaphore another queue waits on). + static constexpr std::array ReleaseSlots{24, 28, 52}; + struct State { Descriptor descriptor{}; std::int32_t dependencies{}; diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h index 11b2f631a5..930e5626aa 100644 --- a/src/video_core/query_cache/query_cache.h +++ b/src/video_core/query_cache/query_cache.h @@ -234,7 +234,12 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type QueryPropertiesFlags flags, u32 payload, u32 subreport) { const bool has_timestamp = True(flags & QueryPropertiesFlags::HasTimeout); const bool is_fence = True(flags & QueryPropertiesFlags::IsAFence); - if (static const bool trace = std::getenv("EDEN_NX2_QUERY_TRACE") != nullptr; trace) { + static const GPUVAddr watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? GPUVAddr{std::strtoull(value, nullptr, 16)} : GPUVAddr{}; + }(); + if (static const bool trace = std::getenv("EDEN_NX2_QUERY_TRACE") != nullptr; + trace || (watch != 0 && (addr & ~GPUVAddr{0xf}) == watch)) { LOG_INFO(HW_GPU, "nx2 report addr={:#x} type={} flags={:#x} payload={:#x}", addr, static_cast(counter_type), static_cast(flags), payload); } diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index 1d9965a9a5..044c783a99 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -1103,6 +1103,14 @@ void RasterizerVulkan::AccelerateInlineToMemory(GPUVAddr address, size_t copy_si if (memory.size() >= sizeof(first_word)) std::memcpy(&first_word, memory.data(), sizeof(first_word)); VideoCommon::TraceNativeQuery(11, gpu_memory->GetID(), address, cpu_addr.value_or(0), 0, first_word, static_cast(copy_size)); + static const GPUVAddr watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? GPUVAddr{std::strtoull(value, nullptr, 16)} : GPUVAddr{}; + }(); + if (watch != 0 && address < watch + 16 && address + copy_size > watch) { + LOG_INFO(Render_Vulkan, "Semaphore watch: inline write {:#x} size {} first word {:#x}", + address, copy_size, first_word); + } } if (!cpu_addr) [[unlikely]] { gpu_memory->WriteBlock(address, memory.data(), copy_size); diff --git a/src/yuzu_cmd/yuzu.cpp b/src/yuzu_cmd/yuzu.cpp index 3aaba78fd1..b46534ec9d 100644 --- a/src/yuzu_cmd/yuzu.cpp +++ b/src/yuzu_cmd/yuzu.cpp @@ -4,6 +4,9 @@ // SPDX-FileCopyrightText: 2014 Citra Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include +#include +#include #include #include #include @@ -34,6 +37,7 @@ #include "core/hle/service/filesystem/filesystem.h" #include "core/loader/loader.h" #include "frontend_common/config.h" +#include "input_common/drivers/keyboard.h" #include "input_common/main.h" #include "network/network.h" #include "sdl_config.h" @@ -327,8 +331,50 @@ extern "C" SDL_AppResult SDL_AppInit(void **appstate, int argc, char **argv) { state->system.InitializeDebugger(); return SDL_APP_CONTINUE; } +// EDEN_NX2_KEY_SCRIPT="seconds:code,seconds:-code,...": press (code) or release (-code) keyboard +// keys, by SDL scancode, that many seconds after start. Lets hidden test runs drive menus. +static void RunKeyScript(InputCommon::InputSubsystem& input_subsystem) { + struct KeyEvent { + double seconds; + int code; + }; + static const std::vector script = [] { + std::vector events; + if (const char* text = std::getenv("EDEN_NX2_KEY_SCRIPT")) { + std::string list{text}; + size_t start = 0; + while (start < list.size()) { + const size_t end = std::min(list.find(',', start), list.size()); + const std::string item = list.substr(start, end - start); + if (const size_t colon = item.find(':'); colon != std::string::npos) { + events.push_back({std::stod(item.substr(0, colon)), + std::stoi(item.substr(colon + 1))}); + } + start = end + 1; + } + } + return events; + }(); + static const auto begin = std::chrono::steady_clock::now(); + static size_t next = 0; + const double elapsed = + std::chrono::duration(std::chrono::steady_clock::now() - begin).count(); + while (next < script.size() && script[next].seconds <= elapsed) { + const int code = script[next].code; + LOG_INFO(Frontend, "Key script: {} key {} at {:.1f} s", code < 0 ? "release" : "press", + code < 0 ? -code : code, elapsed); + if (code < 0) { + input_subsystem.GetKeyboard()->ReleaseKey(static_cast(-code)); + } else { + input_subsystem.GetKeyboard()->PressKey(static_cast(code)); + } + ++next; + } +} + extern "C" SDL_AppResult SDL_AppIterate(void *appstate) { SdlState *state = (SdlState *)appstate; + RunKeyScript(state->input_subsystem); return state->emu_window && state->emu_window->IsOpen() ? SDL_APP_CONTINUE : SDL_APP_SUCCESS; } extern "C" SDL_AppResult SDL_AppEvent(void *appstate, SDL_Event *event) { -- cgit v1.2.3 From e6e0304303ba5280c318e6f66e9d5ea30d375060 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 00:50:13 -0400 Subject: nx2: title 4 runs past its loading screen without GPU stalls - Native QMD: a launch whose shader Eden can't translate still completes and fires its releases (the console never fails a launch; queues wait on them). - QMD releases with flag bit 19 are max releases: NVN2 tracks QMD completion in a word other releases also advance, so a plain write moved it backwards. - Ampere lowering: LDSM.16.M88.{1,2,4} (SHFL + LDS), IMAD with a negated uniform addend, strong LDG (bit 77), 2D TLD4 gather, RED MIN/MAX/..., AST of RZ. 210 of the 214 shaders title 4 uses after its menu now lower. - Maxwell3D: ignore constant buffer binds to slots >= 18 instead of writing past the state array. - Debug: EDEN_NX2_SEM_WATCH covers QMD and compute-macro releases and logs address spaces; EDEN_NX2_3D_METHODS logs the first write to each 3D method; constant buffer binds are logged once per stage/slot. --- .../frontend/ampere/lower_to_maxwell.h | 88 ++++++++++++++++++++-- src/video_core/engines/kepler_compute.cpp | 27 ++++++- src/video_core/engines/maxwell_3d.cpp | 20 +++++ src/video_core/engines/maxwell_3d.h | 2 + src/video_core/engines/puller.cpp | 6 +- src/video_core/native_qmd.h | 9 ++- src/video_core/query_cache/query_cache.h | 5 +- 7 files changed, 141 insertions(+), 16 deletions(-) (limited to 'src/video_core/query_cache/query_cache.h') diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 01dc1d8425..fc3f088c57 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -382,9 +382,18 @@ inline LoweredProgram LowerToMaxwell(std::span code) { break; } case 0xc24: case 0xe24: { // IMAD low word with a uniform multiplier or addend. - if ((insn.words[2] & ~0x080002ffU) != 0x078e0000 || (insn.words[1] & ~0xffU)) + // 0xe24 keeps the uniform addend in bits 32-37; bit 63 negates it ("RZ, RZ, -UR4"). + const bool negate_c = opcode == 0xe24 && insn.Bits(63, 1) != 0; + if ((insn.words[2] & ~0x080002ffU) != 0x078e0000 || + (insn.words[1] & ~(0xffU | (negate_c ? 0x80000000U : 0U)))) return fail(index, "native uniform-addend multiply mode"); - const u32 a = operand(0, false), b = operand(1, false), c = operand(2, false); + const u32 a = operand(0, false), b = operand(1, false); + u32 c = operand(2, false); + if (negate_c) { + emit(0x5c10000000000000ULL | predicate | 252 | (u64{255} << 8) | (u64{c} << 20) | + (u64{1} << 48)); + c = 252; + } multiply_low(a, b, c, false); break; } @@ -654,11 +663,11 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{1} << 29) | (u64{insn.Bits(40, 13)} << 34) | (u64{insn.Bits(81, 3)} << 48)); break; } - case 0x98e: { // RED.E.ADD.STRONG.GPU U32, direct register address. - if (insn.words[2] != 0x0010e100 || insn.Bits(40, 24) || (insn.Bits(24, 8) & 1)) + case 0x98e: { // RED.E.{ADD..XOR}.STRONG.GPU U32 (op in bits 87-89), register address. + if ((insn.words[2] & ~0x03800000U) != 0x0010e100 || insn.Bits(40, 24) || (insn.Bits(24, 8) & 1)) return fail(index, "native global reduction mode"); emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | - (u64{insn.Bits(24, 8)} << 8) | (u64{1} << 48)); + (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(87, 3)} << 23) | (u64{1} << 48)); break; } case 0x2bd: // UFLO.U32, the highest set-bit index (.SH, bit 74: its shift amount). @@ -843,6 +852,37 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{width} << 48)); break; } + case 0x83b: { // LDSM.16.M88.{1,2,4} (ldmatrix): 8x8 16-bit matrices from shared memory. + // Lanes 8m..8m+7 give the row addresses of matrix m; lane t receives the 32 bits at + // row t/4, byte (t%4)*4 of each matrix in Rd+m. Emulated with SHFL.IDX + LDS.32. + const u32 count = 1U << insn.Bits(72, 2); + if (insn.Bits(72, 2) == 3 || (insn.words[2] & ~0x300U) || insn.Bits(32, 8) || + dest + count > 246) + return fail(index, "native matrix load mode"); + u32 address = insn.Bits(24, 8); + const u64 offset = insn.Bits(40, 24); + if (address != 255 && address >= dest && address < dest + count) { + emit(0x5c10000000000000ULL | predicate | 249 | (u64{address} << 8) | (u64{255} << 20)); + address = 249; + } + emit(0xf0c8000000000000ULL | predicate | 250); // S2R laneid + emit(0x3828000000000000ULL | predicate | 251 | (u64{250} << 8) | (u64{2} << 20)); // row + emit(0x0400000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{3} << 20)); + emit(0x3848000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{2} << 20)); // column bytes + for (u32 matrix = 0; matrix < count; ++matrix) { + u32 lane = 251; + if (matrix) { + emit(0x1c00000000000000ULL | predicate | 252 | (u64{251} << 8) | (u64{matrix * 8} << 20)); + lane = 252; + } + emit(0xef10000000000000ULL | predicate | 253 | (u64{address} << 8) | (u64{lane} << 20) | + (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5c10000000000000ULL | predicate | 253 | (u64{253} << 8) | (u64{250} << 20)); + emit(0xef48000000000000ULL | predicate | (dest + matrix) | (u64{253} << 8) | + (offset << 20) | (u64{4} << 48)); + } + break; + } case 0x919: case 0x9c3: { // S2R / S2UR thread and workgroup identities. const u32 special = insn.Bits(72, 8); if (insn.words[1] || insn.words[2] != (special << 8) || @@ -904,7 +944,8 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (size == 5 && (dest & 1)) || (size == 6 && (dest & 3))) return fail(index, "native global load mode"); if (opcode == 0x381) { - if ((insn.words[2] & ~0xe00U) != 0x001e8100 || (address & 1) || address >= 246) + // Bit 77 (strong/scoped ordering) needs nothing extra from a Maxwell LDG. + if ((insn.words[2] & ~0x2e00U) != 0x001e8100 || (address & 1) || address >= 246) return fail(index, "native global load mode"); } else { // The uniform base sits where a register source would (bits 32-37). @@ -1069,6 +1110,33 @@ inline LoweredProgram LowerToMaxwell(std::span code) { return fail(index, "native dependency barrier mode"); } break; + case 0xb63: { // TLD4 (gather), 2D without offsets or depth compare; component at bits 87-88. + const u32 mask = insn.Bits(72, 4), second_dest = insn.Bits(64, 8); + const u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); + const u32 binding = insn.Bits(54, 5); + if (insn.Bits(61, 3) != 1 || insn.Bits(76, 4) || insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || binding >= 18 || insn.Bits(89, 1) || mask == 0 || + (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding)) + return fail(index, "native texture gather mode"); + result.required_texture_buffer = binding; + // Two coordinates: x in Ra, y in Rb (the coordinate list splits evenly across them). + emit(0x5c98000000000000ULL | predicate | 250 | (u64{15} << 39) | (u64{source_a} << 20)); + emit(0x5c98000000000000ULL | predicate | 251 | (u64{15} << 39) | + (u64{source_b == 255 ? 255U : source_b} << 20)); + emit(0xc800000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{255} << 20) | + (u64{2} << 28) | (u64{mask} << 31) | (u64{insn.Bits(40, 14)} << 36) | + (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51) | (u64{insn.Bits(87, 2)} << 56)); + const u32 components = static_cast(std::popcount(mask)); + const bool split = second_dest != 255; + for (u32 i = 0; i < components; ++i) { + u32 target = 255; + if (!split || i < 2) target = dest == 255 ? 255U : static_cast(dest) + i; + else target = second_dest + (i - 2); + if (target != 255) + emit(0x5c98000000000000ULL | predicate | target | (u64{15} << 39) | (u64{250 + i} << 20)); + } + break; + } case 0xb60: { // TEX, scalar 2D implicit-LOD sample const u32 mask = insn.Bits(72, 4); const u32 second_dest = insn.Bits(64, 8); @@ -1226,7 +1294,13 @@ inline LoweredProgram LowerToMaxwell(std::span code) { if (address % 4 || address + components * 4 > 0x400 || insn.Bits(24, 8) != 255) { return fail(index, "native attribute address or indexing"); } - const u32 value_reg = opcode == 0x321 ? static_cast(dest) : insn.Bits(32, 8); + u32 value_reg = opcode == 0x321 ? static_cast(dest) : insn.Bits(32, 8); + if (value_reg == 255) { + if (opcode == 0x321) break; // load into RZ: nothing to do + // AST of RZ writes zeros: store them from zeroed scratch registers. + for (u32 i = 0; i < components; ++i) mov_imm(250 + i, 0); + value_reg = 250; + } if (value_reg + components > 255) return fail(index, "native attribute register range"); const u32 vertex = opcode == 0x321 ? insn.Bits(32, 8) : insn.Bits(64, 8); const u64 translated = (opcode == 0x321 ? 0xefd8000000000000ULL : 0xeff0000000000000ULL) | diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 689f85a1d3..f49dc26f82 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -121,7 +121,8 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar [&](u32 address) { return address < native_registers.size() ? native_registers[address] : 0U; }, [&](u32 address, u32 value) { static u32 diagnostic_writes{}; - if (watched) LOG_INFO(HW_GPU, "Semaphore watch: compute macro emit {:#x} = {:#x}", address, value); + static const bool log_emits = std::getenv("EDEN_NX2_SEM_WATCH_EMITS") != nullptr; + if (watched && log_emits) LOG_INFO(HW_GPU, "Semaphore watch: compute macro emit {:#x} = {:#x}", address, value); if (diagnostic_writes++ < 4096) LOG_DEBUG(HW_GPU, "Native compute MME emit method={:#x}, value={:#x}", address, value); if (address >= 0xe00) return false; CallMethod(system, address, value, true); @@ -154,7 +155,7 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar const GPUVAddr program = (GPUVAddr{native_registers[0xc8 + 49] & 0x1ffff} << 32) | native_registers[0xc8 + 48]; const u32 control = native_registers[0xa7]; - LOG_WARNING(HW_GPU, "Native inline QMD: address={:#x}, program={:#x}, control={:#x}, grid={}/{}/{}", + LOG_DEBUG(HW_GPU, "Native inline QMD: address={:#x}, program={:#x}, control={:#x}, grid={}/{}/{}", address, program, control, native_registers[0xc8 + 12], native_registers[0xc8 + 13] & 0xffff, native_registers[0xc8 + 14] & 0xffff); if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { @@ -236,10 +237,32 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar rasterizer->SignalFence([this, address, payload] { const u32 value = memory_manager.Read(address) + payload; memory_manager.Write(address, value); + static const u64 watch = [] { + const char* text = std::getenv("EDEN_NX2_SEM_WATCH"); + return text ? u64{std::strtoull(text, nullptr, 16)} : u64{}; + }(); + if (watch != 0 && (address & ~u64{0xf}) == watch) { + LOG_INFO(HW_GPU, "Semaphore watch: QMD release add {:#x} -> {:#x}", address, value); + } VideoCommon::TraceNativeQuery(1, memory_manager.GetID(), address, memory_manager.GpuToCpuAddress(address).value_or(0), 0, value, 1U << 26); }); + } else if (flags & (1U << 19)) { + // NVN2 tracks QMD completion with a max release (bit 19) into a word other + // releases also advance: a plain write would move it backwards and leave + // GEQ acquires waiting forever. + rasterizer->SignalFence([this, address, payload] { + const u32 value = std::max(memory_manager.Read(address), payload); + memory_manager.Write(address, value); + }); } else { + static const u64 watch = [] { + const char* text = std::getenv("EDEN_NX2_SEM_WATCH"); + return text ? u64{std::strtoull(text, nullptr, 16)} : u64{}; + }(); + if (watch != 0 && (address & ~u64{0xf}) == (watch & ~u64{0xf})) { + LOG_INFO(HW_GPU, "Semaphore watch: QMD release {:#x} = {:#x} flags {:#x}", address, payload, flags); + } rasterizer->Query(address, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0); } diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp index f86b8e9351..4ef2bbc11b 100644 --- a/src/video_core/engines/maxwell_3d.cpp +++ b/src/video_core/engines/maxwell_3d.cpp @@ -392,6 +392,16 @@ void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, const std::vec draw_manager.DrawDeferred(*this); } +// EDEN_NX2_3D_METHODS: log the first write to each 3D register (finds methods Eden doesn't model). +void Maxwell3D::TraceFirstMethodWrite(u32 method, u32 argument) { + static const bool trace = std::getenv("EDEN_NX2_3D_METHODS") != nullptr; + static std::array seen{}; + if (trace && method < seen.size() && !seen[method]) { + seen[method] = true; + LOG_INFO(HW_GPU, "3D method first write: {:#x} (byte {:#x}) = {:#x}", method, method * 4, argument); + } +} + void Maxwell3D::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { // It is an error to write to a register other than the current macro's ARG register before // it has finished execution. @@ -405,6 +415,7 @@ void Maxwell3D::CallMethod(Core::System& system, u32 method, u32 method_argument ProcessMacro(system, method, &method_argument, 1, is_last_call); return; } + TraceFirstMethodWrite(method, method_argument); ASSERT(method < Regs::NUM_REGS && "Invalid Maxwell3D register, increase the size of the Regs structure"); const u32 argument = ProcessShadowRam(method, method_argument); @@ -605,6 +616,15 @@ void Maxwell3D::ProcessCBBind(size_t stage_index) { // Bind the buffer currently in CB_ADDRESS to the specified index in the desired shader // stage. const auto& bind_data = regs.bind_groups[stage_index]; + static std::array, 5> seen{}; + if (!seen[stage_index][bind_data.shader_slot]) { + seen[stage_index][bind_data.shader_slot] = true; + LOG_INFO(HW_GPU, "Constant buffer bind: stage {} slot {} address {:#x} size {:#x}", stage_index, + bind_data.shader_slot.Value(), regs.const_buffer.Address(), regs.const_buffer.size); + } + if (bind_data.shader_slot >= Regs::MaxConstBuffers) { + return; // slot field is 5 bits; the state array holds 18 + } auto& buffer = state.shader_stages[stage_index].const_buffers[bind_data.shader_slot]; buffer.enabled = bind_data.valid.Value() != 0; buffer.address = regs.const_buffer.Address(); diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h index 29db8022e5..2f67653264 100644 --- a/src/video_core/engines/maxwell_3d.h +++ b/src/video_core/engines/maxwell_3d.h @@ -3135,6 +3135,8 @@ public: /// Reads a register value located at the input method address u32 GetRegisterValue(u32 method) const; + void TraceFirstMethodWrite(u32 method, u32 argument); + /// GPU memory, for Turing macros that read memory into their FIFO or data RAM. MemoryManager& GetMemoryManager() { return memory_manager; diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 047e383b85..f4aaed6036 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -170,8 +170,10 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { return value ? GPUVAddr{std::strtoull(value, nullptr, 16)} : GPUVAddr{}; }(); if (watch != 0 && (regs.semaphore_address.SemaphoreAddress() & ~GPUVAddr{0xf}) == watch) { - LOG_INFO(HW_GPU, "Semaphore watch: channel {} op {} {:#x} value {:#x} (now {:#x})", - dma_pusher.channel_state.bind_id, static_cast(op), + LOG_INFO(HW_GPU, "Semaphore watch: channel {} as {} device {:#x} op {} {:#x} value {:#x} (now {:#x})", + dma_pusher.channel_state.bind_id, dma_pusher.memory_manager.GetID(), + dma_pusher.memory_manager.GpuToCpuAddress(regs.semaphore_address.SemaphoreAddress()).value_or(0), + static_cast(op), regs.semaphore_address.SemaphoreAddress(), regs.semaphore_sequence, dma_pusher.memory_manager.Read(regs.semaphore_address.SemaphoreAddress())); } diff --git a/src/video_core/native_qmd.h b/src/video_core/native_qmd.h index c06f14e6b2..41961c1aea 100644 --- a/src/video_core/native_qmd.h +++ b/src/video_core/native_qmd.h @@ -23,6 +23,7 @@ public: template std::string_view Signal(Address address, Word argument, Read read, Release release, Dispatch dispatch) { struct Pending { Address address; Word action; Word count; }; + std::string_view error; std::array pending{}; std::size_t count = 1; pending[0] = {address, argument & 15, ((argument >> 14) & 0x3ff) + 1}; @@ -73,8 +74,10 @@ public: return "unsupported QMD reduction callback"; } } - if (qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff) && !dispatch(qmd)) - return "native shader dispatch required"; + // A program Eden can't run is skipped, but the QMD still completes: its releases + // are what other queues and the CPU wait on (the console never fails a launch). + const bool dispatched = !(qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff)) || dispatch(qmd); + if (!dispatched) error = "native shader dispatch required"; state.completed = true; for (const Word base : ReleaseSlots) { const Word flags = qmd[base + 1]; @@ -96,7 +99,7 @@ public: } if (qmd[16] & (1U << 5)) return "second dependent QMD unsupported"; } - return {}; + return error; } private: // Release slots (address low, flags | address high, payload). NVN2 patches a third one at diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h index 930e5626aa..f94407b026 100644 --- a/src/video_core/query_cache/query_cache.h +++ b/src/video_core/query_cache/query_cache.h @@ -240,8 +240,9 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type }(); if (static const bool trace = std::getenv("EDEN_NX2_QUERY_TRACE") != nullptr; trace || (watch != 0 && (addr & ~GPUVAddr{0xf}) == watch)) { - LOG_INFO(HW_GPU, "nx2 report addr={:#x} type={} flags={:#x} payload={:#x}", addr, - static_cast(counter_type), static_cast(flags), payload); + LOG_INFO(HW_GPU, "nx2 report addr={:#x} type={} flags={:#x} payload={:#x} as={} device={:#x}", addr, + static_cast(counter_type), static_cast(flags), payload, gpu_memory->GetID(), + gpu_memory->GpuToCpuAddress(addr).value_or(0)); } size_t streamer_id = static_cast(counter_type); auto* streamer = impl->streamers[streamer_id]; -- cgit v1.2.3 From cff4add442d78e088d13005c4c6fdd436b47eecd Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 05:41:45 -0400 Subject: nx2: semaphore releases latch the acquires they satisfy Switch 2 channels that wait for an exact semaphore value poll on the console and see every release. Eden resumes a waiting channel only between command lists and writes releases when fences signal, often several back to back, so an acquire-equal could miss its value and wait forever. Title 4's 3D queue ticks a counter that two other queues wait on value by value; under the loading after the post-customization save one of them missed a tick and the game froze (channel 4 waiting for 0x5293 with the counter at 0x5310). Pending acquires now register with Nx2SemaphoreWaiters before fences run, and every query-cache payload write and QMD add/max release latches the waiters it meets at the moment it is written. --- src/video_core/engines/kepler_compute.cpp | 7 +++ src/video_core/engines/puller.cpp | 20 +++++++- src/video_core/nx2.h | 83 +++++++++++++++++++++++++++++++ src/video_core/query_cache/query_cache.h | 5 +- 4 files changed, 113 insertions(+), 2 deletions(-) (limited to 'src/video_core/query_cache/query_cache.h') diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 85b1f971f7..dc3f9b25b6 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -14,6 +14,7 @@ #include "common/logging.h" #include "shader_recompiler/exception.h" #include "core/core.h" +#include "video_core/nx2.h" #include "video_core/engines/fermi_2d.h" #include "video_core/engines/kepler_compute.h" #include "video_core/engines/maxwell_3d.h" @@ -240,6 +241,9 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar rasterizer->SignalFence([this, address, payload] { const u32 value = memory_manager.Read(address) + payload; memory_manager.Write(address, value); + if (const auto device = memory_manager.GpuToCpuAddress(address)) { + Tegra::Nx2SemaphoreWaiters::Get().OnWrite(*device, value); + } static const u64 watch = [] { const char* text = std::getenv("EDEN_NX2_SEM_WATCH"); return text ? u64{std::strtoull(text, nullptr, 16)} : u64{}; @@ -257,6 +261,9 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar rasterizer->SignalFence([this, address, payload] { const u32 value = std::max(memory_manager.Read(address), payload); memory_manager.Write(address, value); + if (const auto device = memory_manager.GpuToCpuAddress(address)) { + Tegra::Nx2SemaphoreWaiters::Get().OnWrite(*device, value); + } }); } else { static const u64 watch = [] { diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 54abf790d2..b607f20161 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -62,6 +62,15 @@ bool Puller::CanResumeAcquire(DmaPusher& dma_pusher) { return true; } const u32 current = ReadSemaphoreWord(dma_pusher, pending_acquire->address); + if (Nx2SemaphoreWaiters::Get().Latched(this)) { + if (static std::atomic rescued{}; + !pending_acquire->Matches(current) && rescued.fetch_add(1) < 64) { + LOG_INFO(HW_GPU, "Channel {} met semaphore {:#x} value {:#x} at its release (now {:#x})", + dma_pusher.channel_state.bind_id, pending_acquire->address, + pending_acquire->value, current); + } + return true; + } if (pending_acquire->Matches(current)) { return true; } @@ -87,6 +96,7 @@ bool Puller::ResumeAcquire(DmaPusher& dma_pusher) { return false; } pending_acquire.reset(); + Nx2SemaphoreWaiters::Get().Remove(this); return true; } @@ -210,9 +220,17 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { if (acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { return; } + // Registered before the fences run, so a release among them that meets the value counts + // even if a later one overwrites it. + auto& waiters = Nx2SemaphoreWaiters::Get(); + if (const auto device = dma_pusher.memory_manager.GpuToCpuAddress(acquire.address)) { + waiters.Add(this, *device, acquire.value, acquire.operation); + } dma_pusher.rasterizer->ReleaseFences(); - if (!acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { + if (!waiters.Latched(this) && !acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { pending_acquire = acquire; + } else { + waiters.Remove(this); } } else { do { diff --git a/src/video_core/nx2.h b/src/video_core/nx2.h index 7bbbf14785..5b6d89abdd 100644 --- a/src/video_core/nx2.h +++ b/src/video_core/nx2.h @@ -3,8 +3,11 @@ #pragma once +#include #include #include +#include +#include #include "common/common_types.h" @@ -57,4 +60,84 @@ inline u64 Nx2SemaphoreMirror(u64 address) { return 0; } +/// Channels suspended on a semaphore acquire. A console channel polls and sees every release; +/// Eden writes releases in bursts when fences signal and resumes waiters only between command +/// lists, so an acquire-equal could miss its value. Writes latch the waiters they satisfy. +class Nx2SemaphoreWaiters { +public: + // Puller GpuSemaphoreOperation values. + static constexpr u32 AcquireEqual = 0x1; + static constexpr u32 AcquireGequal = 0x4; + static constexpr u32 AcquireMask = 0x8; + + static Nx2SemaphoreWaiters& Get() { + static Nx2SemaphoreWaiters waiters; + return waiters; + } + + void Add(const void* owner, u64 device_address, u32 value, u32 operation) { + std::scoped_lock lk{mutex}; + std::erase_if(waiters, [owner](const Waiter& w) { return w.owner == owner; }); + waiters.push_back({owner, device_address, value, operation, false}); + count.store(waiters.size(), std::memory_order_release); + } + + void Remove(const void* owner) { + if (count.load(std::memory_order_acquire) == 0) { + return; + } + std::scoped_lock lk{mutex}; + std::erase_if(waiters, [owner](const Waiter& w) { return w.owner == owner; }); + count.store(waiters.size(), std::memory_order_release); + } + + [[nodiscard]] bool Latched(const void* owner) { + if (count.load(std::memory_order_acquire) == 0) { + return false; + } + std::scoped_lock lk{mutex}; + return std::ranges::any_of(waiters, + [owner](const Waiter& w) { return w.owner == owner && w.latched; }); + } + + /// A GPU write of `value` to the u32 at `device_address`. + void OnWrite(u64 device_address, u32 value) { + if (count.load(std::memory_order_acquire) == 0) { + return; + } + std::scoped_lock lk{mutex}; + for (auto& w : waiters) { + if (w.address == device_address && !w.latched && Matches(w, value)) { + w.latched = true; + } + } + } + +private: + struct Waiter { + const void* owner; + u64 address; + u32 value; + u32 operation; + bool latched; + }; + + static bool Matches(const Waiter& w, u32 current) { + switch (w.operation) { + case AcquireEqual: + return current == w.value; + case AcquireGequal: + return current >= w.value; + case AcquireMask: + return (current & w.value) != 0; + default: + return false; + } + } + + std::mutex mutex; + std::vector waiters; + std::atomic count{}; +}; + } // namespace Tegra diff --git a/src/video_core/query_cache/query_cache.h b/src/video_core/query_cache/query_cache.h index f94407b026..245ab1996c 100644 --- a/src/video_core/query_cache/query_cache.h +++ b/src/video_core/query_cache/query_cache.h @@ -23,6 +23,7 @@ #include "video_core/gpu.h" #include "video_core/host1x/gpu_device_memory_manager.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" #include "video_core/query_cache/bank_base.h" #include "video_core/query_cache/query_base.h" #include "video_core/query_cache/query_cache_base.h" @@ -273,7 +274,7 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type u8* pointer_timestamp = impl->device_memory.template GetPointer(cpu_addr + 8); bool is_synced = (Settings::IsGPUFenceBehaviorDefault() ? !Settings::IsGPULevelHigh() : !Settings::IsGPUFenceBehaviorBalanced() && !Settings::IsGPUFenceBehaviorAccurate()) && is_fence; std::function operation([this, is_synced, streamer, query_base = query, query_location, - pointer, pointer_timestamp] { + pointer, pointer_timestamp, cpu_addr] { if (True(query_base->flags & QueryFlagBits::IsInvalidated)) { if (!is_synced) [[likely]] { impl->pending_unregister.push_back(query_location); @@ -298,6 +299,7 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type u32 value = static_cast(query_base->value); std::memcpy(pointer, &value, sizeof(value)); } + Tegra::Nx2SemaphoreWaiters::Get().OnWrite(cpu_addr, static_cast(query_base->value)); if (!is_synced) [[likely]] { impl->pending_unregister.push_back(query_location); } @@ -314,6 +316,7 @@ void QueryCacheBase::CounterReport(GPUVAddr addr, QueryType counter_type } else { std::memcpy(pointer, &payload, sizeof(payload)); } + Tegra::Nx2SemaphoreWaiters::Get().OnWrite(cpu_addr, payload); streamer->Free(new_query_id); return; } -- cgit v1.2.3