From 2835fec7fc66e1a1659710de8eddf2dfe67e20f9 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 16:19:35 -0400 Subject: Switch 2 GPU: syncpt server, identity GPU mappings, Ampere classes, submit fences - gsv 204: real syncpoint shared page (values mirrored from host1x, armed waits acknowledged by a 500us poller); first handle is the writable request event. - syncpt 22/201/202: per-slot events signalled through host1x guest actions. - nx2:gpu sub 101 memory is identity-mapped (GPU VA == CPU VA) into every nvhost-as-gpu GMMU via Nx2GpuMemoryRegistry; SMMU physical range widened to 34 bits for the 10/12 GiB layouts. - Puller runs Ampere 3D/compute/copy classes on the Maxwell engines; class 0 unbinds. - Turing+ macros are skipped (Eden only has the Fermi macro ISA). - Submits with increment_value get Eden-side syncpoint increments, since NVN2 signals completion with semaphore releases only. --- src/video_core/engines/puller.cpp | 29 +++++++++++++++++++++++++++++ 1 file changed, 29 insertions(+) (limited to 'src/video_core/engines/puller.cpp') diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 0b05d31718..c9c1153bbf 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -4,6 +4,8 @@ // SPDX-FileCopyrightText: 2022 yuzu Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later +#include + #include "common/assert.h" #include "common/logging.h" #include "common/settings.h" @@ -18,6 +20,7 @@ #include "video_core/engines/puller.h" #include "video_core/gpu.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" #include "video_core/rasterizer_interface.h" namespace Tegra::Engines { @@ -35,6 +38,27 @@ void Puller::ProcessBindMethod(DmaPusher& dma_pusher, const MethodCall& method_c if ((engine & ~Gf100BindClassMask) != 0 && (engine & ~Gf100BindValidMask) == 0) { engine &= Gf100BindClassMask; } + // Switch 2 binds Ampere classes; run them on the Maxwell engines (same method layout for + // most state). Class 0 unbinds the subchannel. + switch (engine & 0xFFFF) { + case 0xC797: // AMPERE_B + engine = static_cast(EngineID::MAXWELL_B); + break; + case 0xC7C0: // AMPERE_COMPUTE_B + engine = static_cast(EngineID::KEPLER_COMPUTE_B); + break; + case 0xC7B5: // AMPERE_DMA_COPY_B + engine = static_cast(EngineID::MAXWELL_DMA_COPY_A); + break; + case 0x0: + if (g_nx2_ampere) { + LOG_DEBUG(HW_GPU, "Switch 2 unbinds subchannel {}", method_call.subchannel); + return; + } + break; + default: + break; + } const auto engine_id = static_cast(engine); bound_engines[method_call.subchannel] = engine_id; @@ -139,6 +163,11 @@ void Puller::ProcessSemaphoreAcquire(DmaPusher& dma_pusher) { /// Calls a GPU puller method. void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { + static const bool nx2_trace = std::getenv("EDEN_NX2_PULLER_TRACE") != nullptr; + if (nx2_trace) { + LOG_INFO(HW_GPU, "nx2 puller sub={} method={:#x} arg={:#x}", method_call.subchannel, + method_call.method, method_call.argument); + } regs.reg_array[method_call.method] = method_call.argument; const auto method = static_cast(method_call.method); -- cgit v1.2.3 From dce3adbc4cae486e2f3c080b995669d6db5ef58f Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 17:04:06 -0400 Subject: Switch 2: window image sizes, queue semaphores, built-in fonts, memory budget - gsv 1510 CalculateRequirements: block-linear size/alignment for window images. - nvhost-gpu: Switch 2 channel event 7 is a real (unsignalled) event. - Puller mirrors NVN2 queue semaphore releases into the page the CPU and GPU wait on (the first two-page nx2:gpu region); EDEN_NX2_METHOD_TRACE logs methods. - pl:u 2000/2010 GetBuiltInFontSharedMemory/GetBuiltInFontData serve plain TTF copies of the shared fonts. - Switch 2 processes count the alias region in their max memory, so the game's nninitStartup heap (total - 6.3 GiB) has room for its sound heap. - gdbrsp.py: interrupt, threads, selthread, hits, continue_async, sleep. --- src/core/file_sys/program_metadata.cpp | 1 + src/core/file_sys/program_metadata.h | 5 ++ src/core/hle/kernel/k_process.cpp | 7 ++ .../hle/service/ns/platform_service_manager.cpp | 56 +++++++++++++++ src/core/hle/service/ns/platform_service_manager.h | 3 + src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp | 13 ++++ src/core/hle/service/nvdrv/devices/nvhost_gpu.h | 2 + src/core/hle/service/nx2_gsv.h | 35 +++++++++ src/video_core/engines/puller.cpp | 20 ++++++ src/video_core/nx2.h | 8 +++ tools/nx2/gdbrsp.py | 83 ++++++++++++++++++++++ 11 files changed, 233 insertions(+) (limited to 'src/video_core/engines/puller.cpp') diff --git a/src/core/file_sys/program_metadata.cpp b/src/core/file_sys/program_metadata.cpp index ad09b19e68..464e496ea3 100644 --- a/src/core/file_sys/program_metadata.cpp +++ b/src/core/file_sys/program_metadata.cpp @@ -47,6 +47,7 @@ Loader::ResultStatus ProgramMetadata::Load(VirtualFile file) { return Loader::ResultStatus::ErrorBadACIDHeader; } LOG_INFO(Loader, "NPDM uses the Switch 2 ACID layout"); + is_switch2 = true; } if (sizeof(AciHeader) != file->ReadObject(&aci_header, npdm_header.aci_offset)) { diff --git a/src/core/file_sys/program_metadata.h b/src/core/file_sys/program_metadata.h index 115e6d6cd1..5ce32af1a4 100644 --- a/src/core/file_sys/program_metadata.h +++ b/src/core/file_sys/program_metadata.h @@ -72,6 +72,10 @@ public: KernelCapabilityDescriptors capabilities); bool Is64BitProgram() const; + /// True when the NPDM uses the Switch 2 ACID layout. + bool IsSwitch2() const { + return is_switch2; + } ProgramAddressSpaceType GetAddressSpaceType() const; u8 GetMainThreadPriority() const; u8 GetMainThreadCore() const; @@ -179,6 +183,7 @@ private: Header npdm_header{}; AciHeader aci_header{}; AcidHeader acid_header{}; + bool is_switch2{}; FileAccessControl acid_file_access{}; FileAccessHeader aci_file_access{}; diff --git a/src/core/hle/kernel/k_process.cpp b/src/core/hle/kernel/k_process.cpp index 27d1b1bfc4..4aac818b77 100644 --- a/src/core/hle/kernel/k_process.cpp +++ b/src/core/hle/kernel/k_process.cpp @@ -1261,6 +1261,13 @@ Result KProcess::LoadFromMetadata(KernelCore& kernel, const FileSys::ProgramMeta // Assign remaining properties. m_ideal_core_id = metadata.GetMainThreadCore(); + // Switch 2 applications get more memory than a 64-bit heap region holds (nninitStartup sizes + // its heap as total memory minus ~6.3 GiB), so count the alias region as well, bounded by the + // pool through the resource limit. + if (metadata.IsSwitch2()) { + m_max_process_memory = m_page_table.GetHeapRegionSize() + m_page_table.GetAliasRegionSize(); + } + // Set up emulation context. this->InitializeInterfaces(kernel); diff --git a/src/core/hle/service/ns/platform_service_manager.cpp b/src/core/hle/service/ns/platform_service_manager.cpp index bf93a2edfa..88aac56b1c 100644 --- a/src/core/hle/service/ns/platform_service_manager.cpp +++ b/src/core/hle/service/ns/platform_service_manager.cpp @@ -22,7 +22,9 @@ #include "core/hle/kernel/k_shared_memory.h" #include "core/hle/kernel/kernel.h" #include "core/hle/service/cmif_serialization.h" +#include "common/alignment.h" #include "core/hle/service/filesystem/filesystem.h" +#include "core/hle/service/ipc_helpers.h" #include "core/hle/service/ns/platform_service_manager.h" namespace Service::NS { @@ -83,6 +85,8 @@ struct IPlatformServiceManager::Impl { // Automatically populated based on shared_fonts dump or system archives. // 6 builtin fonts + extra 2 for whatever may come after boost::container::static_vector shared_font_regions; + /// Plain TTF copies of the same fonts (Switch 2 GetBuiltInFontData), after the BFTTF data. + boost::container::static_vector builtin_font_regions; /// Backing memory for the shared font data std::array shared_font; }; @@ -107,6 +111,8 @@ IPlatformServiceManager::IPlatformServiceManager(Core::System& system_, const ch {106, nullptr, "GetFunctionBlackListVersion"}, {1000, nullptr, "LoadNgWordDataForPlatformRegionChina"}, {1001, nullptr, "GetNgWordDataSizeForPlatformRegionChina"}, + {2000, &IPlatformServiceManager::GetBuiltInFontSharedMemory, "GetBuiltInFontSharedMemory (Switch 2)"}, + {2010, &IPlatformServiceManager::GetBuiltInFontData, "GetBuiltInFontData (Switch 2)"}, }; // clang-format on RegisterHandlers(functions); @@ -116,6 +122,7 @@ IPlatformServiceManager::IPlatformServiceManager(Core::System& system_, const ch // Attempt to load shared font data from disk const auto* nand = fsc.GetSystemNANDContents(); std::size_t offset = 0; + std::vector> plain_fonts; // Rebuild shared fonts from data ncas or synthesize for (auto& font : SHARED_FONTS) { FileSys::VirtualFile romfs; @@ -151,9 +158,58 @@ IPlatformServiceManager::IPlatformServiceManager(Core::System& system_, const ch Common::swap32); // Font offset and size do not account for the header const FontRegion region{u32(offset + 8), u32((font_data_u32.size() * sizeof(u32)) - 8)}; + std::vector plain(region.size); + DecryptSharedFontToTTF(font_data_u32, plain); + plain_fonts.push_back(std::move(plain)); DecryptSharedFont(font_data_u32, impl->shared_font, offset); impl->shared_font_regions.push_back(region); } + + for (const auto& plain : plain_fonts) { + offset = Common::AlignUp(offset, 0x1000); + if (offset + plain.size() > SHARED_FONT_MEM_SIZE) { + LOG_WARNING(Service_NS, "No shared memory left for plain built-in font copies"); + break; + } + std::memcpy(impl->shared_font.data() + offset, plain.data(), plain.size()); + impl->builtin_font_regions.push_back({u32(offset), u32(plain.size())}); + offset += plain.size(); + } +} + +// Switch 2: the shared memory GetBuiltInFontData offsets point into (copy handle, u32 size); the +// SDK maps it read-only. Eden serves its font shared memory, which holds plain copies too. +void IPlatformServiceManager::GetBuiltInFontSharedMemory(HLERequestContext& ctx) { + LOG_WARNING(Service_NS, "(STUBBED) Switch 2 GetBuiltInFontSharedMemory"); + std::memcpy(kernel.GetFontSharedMem().GetPointer(), impl->shared_font.data(), impl->shared_font.size()); + IPC::ResponseBuilder rb{ctx, 3, 1}; + rb.Push(ResultSuccess); + rb.Push(static_cast(SHARED_FONT_MEM_SIZE)); + rb.PushCopyObjects(ctx, kernel.GetFontSharedMem()); +} + +// nn::pl::GetBuiltInFontData(BuiltInFontDataOunce*, BuiltInFontTypeOunce, bool): returns a 16-byte +// bool that says whether the font exists and a 16-byte descriptor (u8 flags, u8 flag bit, u16, +// u32 offset into the font shared memory, u32 size, u32). The client expects plain TrueType/OpenType data +// there. The type numbering is unknown, so types map onto Eden's shared fonts in order. +void IPlatformServiceManager::GetBuiltInFontData(HLERequestContext& ctx) { + IPC::RequestParser rp{ctx}; + const u32 type = rp.Pop(); + const bool flag = rp.Pop() != 0; + const auto& regions = impl->builtin_font_regions; + const bool found = !regions.empty(); + const FontRegion region = found ? regions[std::min(type, regions.size() - 1)] : EMPTY_REGION; + LOG_WARNING(Service_NS, "(STUBBED) Switch 2 GetBuiltInFontData type={} flag={} -> offset={:#x} size={:#x}", + type, flag, region.offset, region.size); + // Raw reply: the bool at offset 0, the descriptor at offset 8. + IPC::ResponseBuilder rb{ctx, 2 + 6}; + rb.Push(ResultSuccess); + rb.Push(found ? 1 : 0); + rb.Push(0); + rb.Push(0); + rb.Push(region.offset); + rb.Push(region.size); + rb.Push(0); } IPlatformServiceManager::~IPlatformServiceManager() = default; diff --git a/src/core/hle/service/ns/platform_service_manager.h b/src/core/hle/service/ns/platform_service_manager.h index 73521abbc2..ed5b1c1e39 100644 --- a/src/core/hle/service/ns/platform_service_manager.h +++ b/src/core/hle/service/ns/platform_service_manager.h @@ -73,6 +73,9 @@ private: Out out_fonts_are_loaded, Out out_font_count, Set::LanguageCode language_code); + void GetBuiltInFontSharedMemory(HLERequestContext& ctx); + void GetBuiltInFontData(HLERequestContext& ctx); + struct Impl; std::unique_ptr impl; }; diff --git a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp index b650d3cac8..a72ccd9b4f 100644 --- a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp +++ b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp @@ -49,6 +49,9 @@ nvhost_gpu::~nvhost_gpu() { events_interface.FreeEvent(sm_exception_breakpoint_int_report_event); events_interface.FreeEvent(sm_exception_breakpoint_pause_report_event); events_interface.FreeEvent(error_notifier_event); + for (auto& [id, event] : nx2_events) { + events_interface.FreeEvent(event); + } syncpoint_manager.FreeSyncpoint(channel_syncpoint); } @@ -543,6 +546,16 @@ Kernel::KEvent* nvhost_gpu::QueryEvent(u32 event_id) { case 3: return error_notifier_event; default: + if (g_nx2_driver) { + // Switch 2 NVN also asks for event 7 and multi-waits on it with its syncpoint events. + // Nothing signals it yet. + auto*& event = nx2_events[event_id]; + if (event == nullptr) { + event = events_interface.CreateEvent(fmt::format("GpuChannelNx2Event{}", event_id)); + } + LOG_WARNING(Service_NVDRV, "(STUBBED) Switch 2 channel event {}", event_id); + return event; + } LOG_CRITICAL(Service_NVDRV, "Unknown Ctrl GPU Event {}", event_id); return nullptr; } diff --git a/src/core/hle/service/nvdrv/devices/nvhost_gpu.h b/src/core/hle/service/nvdrv/devices/nvhost_gpu.h index 70a6b1bd68..ab968d6c7e 100644 --- a/src/core/hle/service/nvdrv/devices/nvhost_gpu.h +++ b/src/core/hle/service/nvdrv/devices/nvhost_gpu.h @@ -6,6 +6,7 @@ #pragma once +#include #include #include #include "common/bit_field.h" @@ -225,6 +226,7 @@ private: Kernel::KEvent* sm_exception_breakpoint_int_report_event; Kernel::KEvent* sm_exception_breakpoint_pause_report_event; Kernel::KEvent* error_notifier_event; + std::map nx2_events; }; } // namespace Service::Nvidia::Devices diff --git a/src/core/hle/service/nx2_gsv.h b/src/core/hle/service/nx2_gsv.h index ef222befca..9d1a4cdc93 100644 --- a/src/core/hle/service/nx2_gsv.h +++ b/src/core/hle/service/nx2_gsv.h @@ -9,6 +9,7 @@ #include #include +#include "common/alignment.h" #include "core/core.h" #include "core/core_timing.h" #include "core/hle/kernel/k_client_session.h" @@ -21,6 +22,7 @@ #include "core/hle/kernel/k_transfer_memory.h" #include "core/memory.h" #include "video_core/host1x/host1x.h" +#include "video_core/nx2.h" #include "core/hle/service/hle_ipc.h" #include "core/hle/service/ipc_helpers.h" #include "core/hle/service/kernel_helpers.h" @@ -291,6 +293,14 @@ private: last, b[6], b[7]); if (last > start) { Nvidia::Devices::Nx2GpuMemoryRegistry::Get().Register(start, last - start + 1); + // The first two-page region holds the queue semaphores: GPU releases go to the second + // page, waits read the first (see video_core/nx2.h). + if (last - start + 1 == 0x2000 && Tegra::g_nx2_semaphore_source_page == 0) { + Tegra::g_nx2_semaphore_mirror_page = start; + Tegra::g_nx2_semaphore_source_page = start + 0x1000; + LOG_WARNING(Service, "nx2:gpu semaphore page {:#x} mirrored to {:#x}", + start + 0x1000, start); + } } IPC::ResponseBuilder rb{ctx, 2}; rb.Push(ResultSuccess); @@ -376,6 +386,7 @@ public: {305, &INx2GraphicsSystemSession::OpenViSession, "OpenViSession (guessed)"}, {497, &INx2GraphicsSystemSession::OpenGpuSession, "OpenGpuSession (guessed)"}, {1402, &INx2GraphicsSystemSession::OpenNvsrvSession, "OpenNvsrvSession (guessed)"}, + {1510, &INx2GraphicsSystemSession::CalculateRequirements, "CalculateRequirements"}, }); } @@ -522,6 +533,30 @@ private: } } + // nn::gsv::shim::CalculateRequirements(u64* size, u64* alignment, preference, u32): the 0x40-byte + // InterModuleUserImagePreference (words 2..17) and a u32 (word 18). A block-linear preference + // (MakeInterModuleBlockUserImagePreference) holds kind 2 at word 3, (width - 1) | (height - 1) + // << 16 at word 4, two u32s at words 6-7 and the pixel format at word 8. Reply: result, u64 + // size, u64 alignment; the client raises the alignment to at least 4 KiB and rounds the size. + void CalculateRequirements(HLERequestContext& ctx) { + const auto* w = ctx.CommandBuffer(); + const u32 width = (w[4] & 0xFFFF) + 1; + const u32 height = (w[4] >> 16) + 1; + const u32 pixel_format = w[8] & 0xFF; + // Pixel format numbering is unknown; 4 bytes per pixel. Oversizing (8 for format 4) left + // the game's 1.5 GiB heap short of its 300 MiB sound heap. + const u64 bytes_per_pixel = 4; + const u64 pitch = Common::AlignUp(width * bytes_per_pixel, 64); + const u64 size = pitch * Common::AlignUp(height, 128); + constexpr u64 alignment = 0x10000; + LOG_WARNING(Service, + "(STUBBED) nx2:gsv TIPC 1510 kind={} {}x{} [6]={:#x} [7]={:#x} format={} " + "[9]={:#x} [18]={:#x} -> size={:#x}", + w[3], width, height, w[6], w[7], pixel_format, w[9], w[18], size); + TipcReplyWords(ctx, {static_cast(size), static_cast(size >> 32), + static_cast(alignment), 0}); + } + void OpenViSession(HLERequestContext& ctx) { // Used by nn::vi2::sdk::Initialize; takes one u64, returns a session (handle at word 3, // result at word 4). diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index c9c1153bbf..9e78d586d1 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -110,6 +110,12 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { const GPUVAddr sequence_address{regs.semaphore_address.SemaphoreAddress()}; const u32 payload = regs.semaphore_sequence; dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); + const u64 source_page = g_nx2_semaphore_source_page; + if (source_page != 0 && (sequence_address & ~GPUVAddr{0xFFF}) == source_page) { + const GPUVAddr mirror = g_nx2_semaphore_mirror_page + (sequence_address & 0xFFF); + LOG_DEBUG(HW_GPU, "Switch 2 semaphore {:#x}={:#x} mirrored to {:#x}", sequence_address, payload, mirror); + dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); + } } else { do { const u32 word = dma_pusher.memory_manager.Read(regs.semaphore_address.SemaphoreAddress()); @@ -234,8 +240,18 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca } /// Calls a GPU engine method. +static bool Nx2MethodTrace() { + static const bool trace = std::getenv("EDEN_NX2_METHOD_TRACE") != nullptr; + return trace; +} + void Puller::CallEngineMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { const EngineID engine = bound_engines[method_call.subchannel]; + if (Nx2MethodTrace()) { + LOG_INFO(HW_GPU, "nx2 method sub={} engine={:04X} method={:#x} arg={:#x}", + method_call.subchannel, static_cast(engine), method_call.method, + method_call.argument); + } switch (engine) { case EngineID::FERMI_TWOD_A: dma_pusher.channel_state.payload->fermi_2d.CallMethod(dma_pusher.system, method_call.method, method_call.argument, method_call.IsLastCall()); @@ -264,6 +280,10 @@ void Puller::CallEngineMethod(DmaPusher& dma_pusher, const MethodCall& method_ca /// Calls a GPU engine multivalue method. void Puller::CallEngineMultiMethod(DmaPusher& dma_pusher, u32 method, u32 subchannel, const u32* base_start, u32 amount, u32 methods_pending) { const EngineID engine = bound_engines[subchannel]; + if (Nx2MethodTrace()) { + LOG_INFO(HW_GPU, "nx2 multimethod sub={} engine={:04X} method={:#x} amount={} first={:#x}", + subchannel, static_cast(engine), method, amount, base_start[0]); + } switch (engine) { case EngineID::FERMI_TWOD_A: dma_pusher.channel_state.payload->fermi_2d.CallMultiMethod(dma_pusher.system, method, base_start, amount, methods_pending); diff --git a/src/video_core/nx2.h b/src/video_core/nx2.h index b8258ca243..9fdf0a7c68 100644 --- a/src/video_core/nx2.h +++ b/src/video_core/nx2.h @@ -5,10 +5,18 @@ #include +#include "common/common_types.h" + namespace Tegra { // Set once a Switch 2 (Ampere, NVN2) driver is talking to the GPU. Its command streams use // Ampere classes and the Turing+ macro ISA, which Eden's Maxwell engines don't implement. inline std::atomic g_nx2_ampere{false}; +// NVN2 queues release their progress values into one GPU page and wait (GPU acquires and CPU +// polls) on the same offsets in the page below it. Something on the console copies one into +// the other; until that is known, the puller mirrors releases that land in the source page. +inline std::atomic g_nx2_semaphore_source_page{0}; // GPU VA of the released-to page +inline std::atomic g_nx2_semaphore_mirror_page{0}; // GPU VA of the waited-on page + } // namespace Tegra diff --git a/tools/nx2/gdbrsp.py b/tools/nx2/gdbrsp.py index 9b3b2d8b37..f0acd2aebd 100644 --- a/tools/nx2/gdbrsp.py +++ b/tools/nx2/gdbrsp.py @@ -131,6 +131,10 @@ def main(): break elif op == "continue": print("stop:", rsp.cmd("c")) + elif op == "continue_async": + # Resume without waiting for a stop reply (pair with interrupt). + pkt = b"c" + rsp.s.sendall(b"$" + pkt + b"#" + f"{sum(pkt) & 0xff:02x}".encode()) elif op == "x0": r = rsp.regs() regs = r @@ -156,6 +160,85 @@ def main(): if nfp <= fp: break fp = nfp + elif op == "interrupt": + # Stop a running target (Ctrl-C) and print the stop reply. + rsp.s.sendall(b"") + print("interrupt:", rsp._read_packet()) + elif op == "threads": + # threads [frames]: every thread's pc, lr and frame-pointer chain. + ids = [] + r = rsp.cmd("qfThreadInfo") + while r.startswith("m"): + ids += [t for t in r[1:].split(",") if t] + r = rsp.cmd("qsThreadInfo") + depth = int(parts[1]) if len(parts) > 1 else 16 + for t in ids: + rsp.cmd(f"Hg{t}") + r = rsp.regs() + chain = [] + fp = r.get("x29", 0) + for _ in range(depth): + d = rsp.mem(fp, 16) + if len(d) < 16: + break + nfp, ret = struct.unpack(" : select the first thread whose pc, lr or frame chain falls in + # [lo, hi); later regs/mem ops then read that thread. + lo, hi = int(parts[1], 16), int(parts[2], 16) + ids = [] + r = rsp.cmd("qfThreadInfo") + while r.startswith("m"): + ids += [t for t in r[1:].split(",") if t] + r = rsp.cmd("qsThreadInfo") + for t in ids: + rsp.cmd(f"Hg{t}") + r = rsp.regs() + cands = [r.get("pc", 0), r.get("x30", 0) & 0xFFFFFFFFFFFF] + fp = r.get("x29", 0) + for _ in range(8): + d = rsp.mem(fp, 16) + if len(d) < 16: + break + nfp, ret = struct.unpack(" [memreg]: stop at addr n times, printing x0/x1/lr (and 16 bytes at + # memreg) each time; steps off the breakpoint before resuming. + addr, n = int(parts[1], 16), int(parts[2]) + memreg = parts[3] if len(parts) > 3 else None + for i in range(n): + rsp.cmd(f"Z0,{addr:x},4") + rsp.cmd("c") + r = rsp.regs() + regs = r + extra = "" + if memreg and memreg in r: + extra = " mem=" + rsp.mem(r[memreg], 16).hex() + print(f" hit {i} pc={r.get('pc', 0):#x} x0={r.get('x0', 0):#x} x1={r.get('x1', 0):#x} " + f"x19={r.get('x19', 0):#x} x20={r.get('x20', 0):#x}" + extra) + if r.get("pc") != addr: + break + rsp.cmd(f"z0,{addr:x},4") + rsp.cmd("s") + elif op == "sleep": + import time + time.sleep(float(parts[1])) elif op == "quit": break rsp.s.close() -- cgit v1.2.3 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/core/hle/service/nvdrv/devices/nvhost_gpu.cpp | 12 ++ src/video_core/control/channel_state.cpp | 6 +- src/video_core/control/scheduler.cpp | 32 +++++ src/video_core/control/scheduler.h | 8 ++ src/video_core/dma_pusher.cpp | 157 ++++++++++++++++++++-- src/video_core/dma_pusher.h | 49 ++++++- src/video_core/engines/kepler_compute.cpp | 15 ++- src/video_core/engines/maxwell_dma.cpp | 6 +- src/video_core/engines/puller.cpp | 133 +++++++++++++++++- src/video_core/engines/puller.h | 23 ++++ src/video_core/gpu_thread.cpp | 14 +- src/video_core/memory_manager.cpp | 5 +- src/video_core/nx2.h | 35 ++++- src/video_core/query_cache/query_cache.h | 5 + 14 files changed, 473 insertions(+), 27 deletions(-) (limited to 'src/video_core/engines/puller.cpp') diff --git a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp index a72ccd9b4f..cf3948c321 100644 --- a/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp +++ b/src/core/hle/service/nvdrv/devices/nvhost_gpu.cpp @@ -4,7 +4,9 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include +#include #include "common/assert.h" #include "common/logging.h" #include "core/core.h" @@ -422,6 +424,16 @@ NvResult nvhost_gpu::SubmitGPFIFOImpl(IoctlSubmitGpfifo& params, Tegra::CommandL LOG_DEBUG(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; + if (check) { + std::vector words; + for (const auto& entry : entries.command_lists) { + words.resize(entry.size); + system.ApplicationMemory().ReadBlock(entry.addr.Value(), words.data(), + words.size() * sizeof(u32)); + entries.submit_hashes.push_back(Tegra::HashPushbufferWords(words)); + } + } } auto& flags = params.flags; diff --git a/src/video_core/control/channel_state.cpp b/src/video_core/control/channel_state.cpp index 12a452dfc0..a1a8110510 100644 --- a/src/video_core/control/channel_state.cpp +++ b/src/video_core/control/channel_state.cpp @@ -21,6 +21,7 @@ namespace { // Match NVK/Nouveau's initial pushbuffer subchannel layout. constexpr u32 Nvk3DSubchannel = 0; constexpr u32 NvkComputeSubchannel = 1; +constexpr u32 NvkM2mfSubchannel = 2; constexpr u32 Nvk2DSubchannel = 3; constexpr u32 NvkCopySubchannel = 4; @@ -29,7 +30,10 @@ void BindNvkDefaultSubchannels(ChannelState::Payload& payload) { dma_pusher.BindSubchannel(&payload.maxwell_3d, Nvk3DSubchannel, Engines::EngineTypes::Maxwell3D); dma_pusher.BindSubchannel(&payload.kepler_compute, NvkComputeSubchannel, Engines::EngineTypes::KeplerCompute); - // Subchannel 2 is M2MF there; Eden does not expose a 0x9039 engine yet. + // Subchannel 2 is M2MF/inline-to-memory there (NVN uses it without binding a class); + // KeplerMemory has the same LINE_LENGTH_IN/OFFSET_OUT/LAUNCH_DMA/LOAD_INLINE_DATA layout. + dma_pusher.BindSubchannel(&payload.kepler_memory, NvkM2mfSubchannel, + Engines::EngineTypes::KeplerMemory); dma_pusher.BindSubchannel(&payload.fermi_2d, Nvk2DSubchannel, Engines::EngineTypes::Fermi2D); dma_pusher.BindSubchannel(&payload.maxwell_dma, NvkCopySubchannel, Engines::EngineTypes::MaxwellDMA); diff --git a/src/video_core/control/scheduler.cpp b/src/video_core/control/scheduler.cpp index 8af13e415c..24f1db1f7c 100644 --- a/src/video_core/control/scheduler.cpp +++ b/src/video_core/control/scheduler.cpp @@ -4,7 +4,9 @@ // SPDX-FileCopyrightText: 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later +#include #include +#include #include "common/assert.h" #include "video_core/control/channel_state.h" @@ -26,6 +28,36 @@ void Scheduler::Push(GPU& gpu, s32 channel, CommandList&& entries) { // Multiple channels can prepare their commands in parallel. channel_state->payload->dma_pusher.Push(std::move(entries)); channel_state->payload->dma_pusher.DispatchCalls(); + if (channel_state->payload->dma_pusher.HasPendingAcquire() && + std::find(pending_channels.begin(), pending_channels.end(), channel_state) == + pending_channels.end()) { + pending_channels.push_back(channel_state); + } + if (!pending_channels.empty()) { + ResumePending(gpu); + } +} + +void Scheduler::ResumePending(GPU& gpu) { + // Running one channel can release the semaphore another waits for: retry until none moves. + for (bool progress = true; progress;) { + progress = false; + for (const auto& state : pending_channels) { + auto& dma_pusher = state->payload->dma_pusher; + if (!dma_pusher.HasPendingAcquire() || !dma_pusher.CanResumeAcquire()) { + continue; + } + { + std::scoped_lock lk{scheduling_guard}; + gpu.BindChannel(state->bind_id); + } + dma_pusher.DispatchCalls(); + progress = true; + } + } + std::erase_if(pending_channels, [](const auto& state) { + return !state->payload->dma_pusher.HasPendingAcquire(); + }); } void Scheduler::DeclareChannel(std::shared_ptr new_channel) { diff --git a/src/video_core/control/scheduler.h b/src/video_core/control/scheduler.h index 342bea5025..3a3dcc58cf 100644 --- a/src/video_core/control/scheduler.h +++ b/src/video_core/control/scheduler.h @@ -8,6 +8,7 @@ #include #include +#include #include "common/container/unordered_map.h" #include "video_core/dma_pusher.h" @@ -24,11 +25,18 @@ class Scheduler { public: void Push(GPU& gpu, s32 channel, CommandList&& entries); + /// Runs channels whose pending semaphore acquire is now met. + void ResumePending(GPU& gpu); + [[nodiscard]] bool HasPendingChannels() const { + return !pending_channels.empty(); + } + void DeclareChannel(std::shared_ptr new_channel); private: ::Common::unordered_map> channels; std::mutex scheduling_guard; + std::vector> pending_channels; ///< Suspended on an acquire }; } // namespace Control diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp index 2c0ba007ef..0889964337 100644 --- a/src/video_core/dma_pusher.cpp +++ b/src/video_core/dma_pusher.cpp @@ -4,13 +4,17 @@ // SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include + #include "common/settings.h" #include "core/core.h" +#include "core/memory.h" #include "video_core/dma_pusher.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/gpu.h" #include "video_core/guest_memory.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" #include "video_core/rasterizer_interface.h" namespace Tegra { @@ -29,7 +33,7 @@ DmaPusher::DmaPusher(Core::System& system_, MemoryManager& memory_manager_, Cont DmaPusher::~DmaPusher() = default; void DmaPusher::DispatchCalls() { - dma_pushbuffer_subindex = 0; + // The sub-index is kept: a channel suspended on a semaphore resumes inside its command list. dma_state.is_last_call = true; while (system.IsPoweredOn()) { if (!Step()) { @@ -44,6 +48,9 @@ bool DmaPusher::Step() { if (!ib_enable || dma_pushbuffer.empty()) { return false; } + if (!puller.ResumeAcquire(*this)) { + return false; + } CommandList& command_list = dma_pushbuffer.front(); @@ -57,7 +64,13 @@ bool DmaPusher::Step() { } if (prefetch_size > 0) { - ProcessCommands(command_list.prefetch_command_list); + const std::span commands{command_list.prefetch_command_list}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); + if (puller.HasPendingAcquire()) { + return false; + } + command_resume_offset = 0; dma_pushbuffer.pop(); return true; } @@ -83,13 +96,33 @@ bool DmaPusher::Step() { subchannels[dma_state.subchannel]->current_dirty = memory_manager.IsMemoryDirty(dma_state.dma_get, words * sizeof(u32)); } } + if (command_resume_offset == 0) { + if (g_nx2_ampere) { + ResolveSplitMethod(header.addr); + } + if (dma_pushbuffer_subindex < command_list.submit_hashes.size()) { + CheckSubmitHash(header, command_list.submit_hashes[dma_pushbuffer_subindex]); + } + } + // A resumed entry is read again: the words after the acquire may be new by now. const bool use_safe = Settings::IsDMALevelDefault() ? Settings::IsGPULevelHigh() : Settings::IsDMALevelSafe(); if (use_safe) { Tegra::Memory::GpuGuestMemoryheaders(memory_manager, dma_state.dma_get, header.size, &command_headers); - ProcessCommands(headers); + const std::span commands{headers.data(), headers.size()}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); } else { Tegra::Memory::GpuGuestMemoryheaders(memory_manager, dma_state.dma_get, header.size, &command_headers); - ProcessCommands(headers); + const std::span commands{headers.data(), headers.size()}; + command_resume_offset += + ProcessCommands(commands.subspan(command_resume_offset), command_resume_offset); + } + if (puller.HasPendingAcquire()) { + return false; + } + command_resume_offset = 0; + if (dma_state.method_count != 0) { + split_method_end = header.addr + header.size * sizeof(u32); } } @@ -111,20 +144,23 @@ bool DmaPusher::Step() { return true; } -void DmaPusher::ProcessCommands(std::span commands) { +size_t DmaPusher::ProcessCommands(std::span commands, size_t word_offset) { for (size_t index = 0; index < commands.size();) { // Data word of methods command if (dma_state.method_count && dma_state.non_incrementing) { auto const& command_header = commands[index]; //must ref (MUltiMethod re) - dma_state.dma_word_offset = u32(index * sizeof(u32)); - const u32 max_write = u32(std::min(index + dma_state.method_count, commands.size()) - index); + dma_state.dma_word_offset = u64((word_offset + index) * sizeof(u32)); + u32 max_write = u32(std::min(index + dma_state.method_count, commands.size()) - index); + if (g_nx2_ampere && dma_state.method < non_puller_methods) { + max_write = 1; // a semaphore acquire may suspend the channel after any word + } CallMultiMethod(&command_header.argument, max_write); dma_state.method_count -= max_write; dma_state.is_last_call = true; index += max_write; } else if (dma_state.method_count) { auto const command_header = commands[index]; //can copy - dma_state.dma_word_offset = u32(index * sizeof(u32)); + dma_state.dma_word_offset = u64((word_offset + index) * sizeof(u32)); dma_state.is_last_call = dma_state.method_count <= 1; CallMethod(command_header.argument); dma_state.method += !dma_state.non_incrementing ? 1 : 0; @@ -163,6 +199,46 @@ void DmaPusher::ProcessCommands(std::span commands) { } index++; } + if (puller.HasPendingAcquire()) { + return index; + } + } + return commands.size(); +} + +// NVN on Switch 2 ends GPFIFO entries in the middle of a method's data and calls another command +// buffer (e.g. a compute dispatch) before the entry that carries the rest, which starts right +// where the first one ended. An entry starting anywhere else is such a call: run it with a clean +// method state and resume the unfinished method when the stream continues at that address. +void DmaPusher::ResolveSplitMethod(GPUVAddr entry_address) { + if (dma_state.method_count != 0 && entry_address != split_method_end) { + if (parked_method) { + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, + "Dropping parked method {:#x} (sub {}, {} words left, resume {:#x})", + parked_method->method, parked_method->subchannel, + parked_method->method_count, parked_method->resume_address); + } + } + parked_method = ParkedMethod{ + .method = dma_state.method, + .subchannel = dma_state.subchannel, + .method_count = dma_state.method_count, + .non_incrementing = dma_state.non_incrementing, + .increment_once = dma_increment_once, + .resume_address = split_method_end, + }; + dma_state.method_count = 0; + dma_increment_once = false; + } else if (dma_state.method_count == 0 && parked_method && + entry_address == parked_method->resume_address) { + dma_state.method = parked_method->method; + dma_state.subchannel = parked_method->subchannel; + dma_state.method_count = parked_method->method_count; + dma_state.non_incrementing = parked_method->non_incrementing; + dma_increment_once = parked_method->increment_once; + parked_method.reset(); } } @@ -182,6 +258,10 @@ void DmaPusher::CallMethod(u32 argument) { }); } else { auto subchannel = subchannels[dma_state.subchannel]; + if (subchannel == nullptr) { + ReportUnboundSubchannel(); + return; + } if (!subchannel->execution_mask[dma_state.method]) { subchannel->method_sink.emplace_back(dma_state.method, argument); } else { @@ -197,12 +277,73 @@ void DmaPusher::CallMultiMethod(const u32* base_start, u32 num_methods) { puller.CallMultiMethod(*this, dma_state.method, dma_state.subchannel, base_start, num_methods, dma_state.method_count); } else { auto subchannel = subchannels[dma_state.subchannel]; + if (subchannel == nullptr) { + ReportUnboundSubchannel(); + return; + } subchannel->ConsumeSink(system); subchannel->current_dma_segment = dma_state.dma_get + dma_state.dma_word_offset; subchannel->CallMultiMethod(system, dma_state.method, base_start, num_methods, dma_state.method_count); } } +void DmaPusher::ReportUnboundSubchannel() { + // Methods sent to a subchannel with no engine (no SetObject yet) are dropped. + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, "Dropping method {:#x} on unbound subchannel {}", dma_state.method, + dma_state.subchannel); + } +} + +void DmaPusher::CheckSubmitHash(const CommandListHeader& header, u64 submit_hash) { + std::vector words(header.size); + memory_manager.ReadBlockUnsafe(header.addr, words.data(), words.size() * sizeof(u32)); + static std::atomic checked{}; + static std::atomic changed{}; + ++checked; + if (HashPushbufferWords(words) != submit_hash) { + if (changed.fetch_add(1) < 16) { + LOG_WARNING(HW_GPU, "pushbuffer entry {:#x}+{:#x} changed between submit and execution", + header.addr.Value(), header.size.Value()); + } + } + if ((checked % 20000) == 0) { + LOG_WARNING(HW_GPU, "pushbuffer check: {} of {} entries changed after submit", + changed.load(), checked.load()); + } +} + +void DmaPusher::DumpCurrentEntry(const char* reason) { + LOG_WARNING(HW_GPU, "pushbuffer dump ({}): method={:#x} sub={} count={} non_inc={} get={:#x}", + reason, dma_state.method, dma_state.subchannel, dma_state.method_count, + dma_state.non_incrementing, dma_state.dma_get); + if (dma_pushbuffer.empty()) { + return; + } + const auto& lists = dma_pushbuffer.front().command_lists; + const size_t first = dma_pushbuffer_subindex >= 2 ? dma_pushbuffer_subindex - 2 : 0; + for (size_t i = first; i <= dma_pushbuffer_subindex && i < lists.size(); ++i) { + const CommandListHeader& header = lists[i]; + LOG_WARNING(HW_GPU, " entry {} of {}: addr={:#x} size={:#x}", i, lists.size(), + header.addr.Value(), header.size.Value()); + for (u64 word = 0; word < std::min(header.size, 0x80); word += 8) { + std::string gpu_line; + std::string cpu_line; + for (u64 j = word; j < std::min(header.size, word + 8); ++j) { + const GPUVAddr address = header.addr + j * sizeof(u32); + const u32 gpu_value = memory_manager.Read(address); + const u32 cpu_value = system.ApplicationMemory().Read32(address); + gpu_line += fmt::format(" {:08x}", gpu_value); + cpu_line += gpu_value == cpu_value ? std::string(" ........") + : fmt::format(" {:08x}", cpu_value); + } + LOG_WARNING(HW_GPU, " +{:03x}:{}", word * 4, gpu_line); + LOG_WARNING(HW_GPU, " cpu :{}", cpu_line); + } + } +} + void DmaPusher::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) { rasterizer = rasterizer_; } diff --git a/src/video_core/dma_pusher.h b/src/video_core/dma_pusher.h index 62af628511..e735839d30 100644 --- a/src/video_core/dma_pusher.h +++ b/src/video_core/dma_pusher.h @@ -8,6 +8,7 @@ #include #include +#include #include #include #include @@ -114,8 +115,19 @@ struct CommandList final { boost::container::small_vector command_lists; boost::container::small_vector prefetch_command_list; + /// EDEN_NX2_PB_CHECK: hash of each entry's words when it was submitted. + std::vector submit_hashes; }; +/// FNV-1a over pushbuffer words (EDEN_NX2_PB_CHECK). +inline u64 HashPushbufferWords(std::span words) { + u64 hash = 0xcbf29ce484222325ULL; + for (const u32 word : words) { + hash = (hash ^ word) * 0x100000001b3ULL; + } + return hash; +} + /// @brief The DmaPusher class implements DMA submission to FIFOs, providing an area of memory that the /// emulated app fills with commands and tells PFIFO to process. The pushbuffers are then assembled /// into a "command stream" consisting of 32-bit words that make up "commands". @@ -132,6 +144,14 @@ public: void DispatchCalls(); + /// Switch 2: true while the channel is suspended on a semaphore acquire. + [[nodiscard]] bool HasPendingAcquire() const { + return puller.HasPendingAcquire(); + } + [[nodiscard]] bool CanResumeAcquire() { + return puller.CanResumeAcquire(*this); + } + void BindSubchannel(Engines::EngineInterface* engine, u32 subchannel_id, Engines::EngineTypes engine_type) { subchannels[subchannel_id] = engine; subchannel_type[subchannel_id] = engine_type; @@ -143,12 +163,24 @@ private: static constexpr u32 non_puller_methods = 0x40; static constexpr u32 max_subchannels = 8; bool Step(); - void ProcessCommands(std::span commands); + /// Returns how many words were processed: fewer than given when a semaphore acquire + /// suspended the channel. + size_t ProcessCommands(std::span commands, size_t word_offset = 0); + size_t command_resume_offset{}; ///< Words of the current entry done before a suspension void SetState(const CommandHeader& command_header); void CallMethod(u32 argument); void CallMultiMethod(const u32* base_start, u32 num_methods); + void ReportUnboundSubchannel(); + +public: + /// Logs the GPFIFO entry being processed (and the one before it), read through the GPU MMU + /// and straight from guest memory, to find pushbuffer desyncs. Debug aid for Switch 2. + void DumpCurrentEntry(const char* reason); + void CheckSubmitHash(const CommandListHeader& header, u64 submit_hash); + +private: public: Common::ScratchBuffer command_headers; ///< Buffer for list of commands fetched at once @@ -172,6 +204,21 @@ public: Control::ChannelState& channel_state; DmaState dma_state{}; + + // Switch 2: a method whose data continues in a later GPFIFO entry, set aside while another + // command buffer called in between runs. + struct ParkedMethod { + u32 method; + u32 subchannel; + u32 method_count; + bool non_incrementing; + bool increment_once; + GPUVAddr resume_address; ///< Where the entry that continues the method starts + }; + void ResolveSplitMethod(GPUVAddr entry_address); + std::optional parked_method; + GPUVAddr split_method_end{}; ///< End of the last entry that ended inside a method + std::array subchannels{}; std::array subchannel_type; diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 96a12539a2..166bbd6d53 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -35,13 +35,24 @@ void KeplerCompute::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) void KeplerCompute::ConsumeSinkImpl(Core::System& system) { for (auto [method, value] : method_sink) { - regs.reg_array[method] = value; + if (method < Regs::NUM_REGS) { + regs.reg_array[method] = value; + } } method_sink.clear(); } void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { - ASSERT_MSG(method < Regs::NUM_REGS, "Invalid KeplerCompute register, increase the size of the Regs structure"); + if (method >= Regs::NUM_REGS) { + // Ampere compute (Switch 2) uses methods past Kepler's register file; writing them would + // run off reg_array, so drop them. + static bool reported = false; + if (!reported) { + reported = true; + LOG_WARNING(HW_GPU, "Ignoring KeplerCompute method {:#x} past the register file", method); + } + return; + } regs.reg_array[method] = method_argument; diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp index ed4675ed12..5ab953e510 100644 --- a/src/video_core/engines/maxwell_dma.cpp +++ b/src/video_core/engines/maxwell_dma.cpp @@ -42,7 +42,11 @@ void MaxwellDMA::ConsumeSinkImpl(Core::System& system) { } void MaxwellDMA::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { - ASSERT_MSG(method < NUM_REGS, "Invalid MaxwellDMA register"); + if (method >= NUM_REGS) { + // Ampere copy (Switch 2) can use methods past this register file; don't write past it. + LOG_DEBUG(HW_GPU, "Ignoring MaxwellDMA method {:#x} past the register file", method); + return; + } regs.reg_array[method] = method_argument; diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index 9e78d586d1..97a58903fe 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: 2022 yuzu Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later +#include #include #include "common/assert.h" @@ -30,8 +31,58 @@ namespace { constexpr u32 Gf100BindClassMask = 0xffff; constexpr u32 Gf100BindValidMask = 0x1f0000 | Gf100BindClassMask; +// Reads a semaphore word, flushing a GPU-written copy out of the buffer cache first. +u32 ReadSemaphoreWord(DmaPusher& dma_pusher, GPUVAddr address) { + if (const auto device = dma_pusher.memory_manager.GpuToCpuAddress(address); + device && dma_pusher.rasterizer->MustFlushRegion(*device, sizeof(u32), + VideoCommon::CacheType::BufferCache)) { + dma_pusher.rasterizer->FlushRegion(*device, sizeof(u32), + VideoCommon::CacheType::BufferCache); + } + return dma_pusher.memory_manager.Read(address); +} + } // Anonymous namespace +bool Puller::PendingAcquire::Matches(u32 current) const { + switch (static_cast(operation)) { + case GpuSemaphoreOperation::AcquireEqual: + return current == value; + case GpuSemaphoreOperation::AcquireGequal: + return current >= value; + case GpuSemaphoreOperation::AcquireMask: + return (current & value) != 0; + default: + return true; + } +} + +bool Puller::CanResumeAcquire(DmaPusher& dma_pusher) { + if (!pending_acquire) { + return true; + } + const u32 current = ReadSemaphoreWord(dma_pusher, pending_acquire->address); + if (pending_acquire->Matches(current)) { + return true; + } + 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", + dma_pusher.channel_state.bind_id, pending_acquire->address, + pending_acquire->operation, pending_acquire->value, current); + } + return false; +} + +bool Puller::ResumeAcquire(DmaPusher& dma_pusher) { + if (!CanResumeAcquire(dma_pusher)) { + return false; + } + pending_acquire.reset(); + return true; +} + void Puller::ProcessBindMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { LOG_DEBUG(HW_GPU, "Binding subchannel {} to engine {:#x}", method_call.subchannel, method_call.argument); u32 engine = method_call.argument; @@ -109,12 +160,39 @@ void Puller::ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher) { if (op == GpuSemaphoreOperation::WriteLong) { const GPUVAddr sequence_address{regs.semaphore_address.SemaphoreAddress()}; const u32 payload = regs.semaphore_sequence; - dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); - const u64 source_page = g_nx2_semaphore_source_page; - if (source_page != 0 && (sequence_address & ~GPUVAddr{0xFFF}) == source_page) { - const GPUVAddr mirror = g_nx2_semaphore_mirror_page + (sequence_address & 0xFFF); + // RELEASE_SIZE (bit 24) set means a 4-byte release without the timestamp; writing 16 + // bytes there clobbers the words next to it (Switch 2 packs semaphores tightly). + const auto properties = g_nx2_ampere && (regs.semaphore_trigger & (1U << 24)) != 0 + ? VideoCommon::QueryPropertiesFlags::IsAFence + : VideoCommon::QueryPropertiesFlags::HasTimeout; + dma_pusher.rasterizer->Query(sequence_address, VideoCommon::QueryType::Payload, properties, payload, 0); + if (const GPUVAddr mirror = g_nx2_ampere ? Nx2SemaphoreMirror(sequence_address) : 0; + mirror != 0) { LOG_DEBUG(HW_GPU, "Switch 2 semaphore {:#x}={:#x} mirrored to {:#x}", sequence_address, payload, mirror); - dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, VideoCommon::QueryPropertiesFlags::HasTimeout, payload, 0); + dma_pusher.rasterizer->Query(mirror, VideoCommon::QueryType::Payload, properties, payload, 0); + } + } else if (g_nx2_ampere) { + // Switch 2 waits on semaphores the CPU (or another channel) releases, e.g. once it has + // written the commands this channel runs next: suspend the channel until the value is + // met instead of reading ahead (approach from changes.patch). + PendingAcquire acquire{ + .address = regs.semaphore_address.SemaphoreAddress(), + .value = regs.semaphore_sequence, + .operation = static_cast(op), + .since = std::chrono::steady_clock::now(), + .reported = false, + }; + if (op != GpuSemaphoreOperation::AcquireEqual && op != GpuSemaphoreOperation::AcquireGequal && + op != GpuSemaphoreOperation::AcquireMask) { + LOG_ERROR(HW_GPU, "Invalid semaphore operation {}", static_cast(op)); + return; + } + if (acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { + return; + } + dma_pusher.rasterizer->ReleaseFences(); + if (!acquire.Matches(ReadSemaphoreWord(dma_pusher, acquire.address))) { + pending_acquire = acquire; } } else { do { @@ -167,6 +245,36 @@ void Puller::ProcessSemaphoreAcquire(DmaPusher& dma_pusher) { } } +// Ampere SEM_EXECUTE (method 0x1B): operation in bits 2:0 on the semaphore at +// SEM_ADDR_HI:SEM_ADDR_LO (methods 0x18/0x17) with payload SEM_PAYLOAD_LO (0x19). +void Puller::ProcessSemaphoreExecute(DmaPusher& dma_pusher, u32 argument) { + const GPUVAddr address = (GPUVAddr{regs.reg_array[0x18] & 0x1FFFFFF} << 32) | + (regs.reg_array[0x17] & ~3U); + const u32 payload = regs.reg_array[0x19]; + const u32 operation = argument & 7; + if (operation == 1) { // RELEASE + dma_pusher.rasterizer->Query(address, VideoCommon::QueryType::Payload, + VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0); + return; + } + // Acquires: check once after flushing pending fences; never block the GPU thread on them. + dma_pusher.rasterizer->ReleaseFences(); + const u32 word = dma_pusher.memory_manager.Read(address); + const bool ready = operation == 0 ? word == payload + : operation == 2 ? word >= payload + : operation == 3 ? static_cast(word - payload) >= 0 + : operation == 4 ? (word & payload) != 0 + : operation == 5 ? (word & payload) == 0 + : true; + if (!ready) { + static std::atomic reported{}; + if (reported.fetch_add(1) < 8) { + LOG_WARNING(HW_GPU, "SEM_EXECUTE op {} at {:#x}: {:#x} vs {:#x} not met, continuing", + operation, address, word, payload); + } + } +} + /// Calls a GPU puller method. void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_call) { static const bool nx2_trace = std::getenv("EDEN_NX2_PULLER_TRACE") != nullptr; @@ -221,11 +329,18 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca break; } case BufferMethods::SemaphoreAcquire: { - ProcessSemaphoreAcquire(dma_pusher); + // On Ampere host classes 0x17..0x1A are SEM_ADDR_LO/HI and SEM_PAYLOAD_LO/HI. + if (!g_nx2_ampere) { + ProcessSemaphoreAcquire(dma_pusher); + } break; } case BufferMethods::SemaphoreRelease: { - ProcessSemaphoreRelease(dma_pusher); + if (g_nx2_ampere) { + ProcessSemaphoreExecute(dma_pusher, method_call.argument); + } else { + ProcessSemaphoreRelease(dma_pusher); + } break; } case BufferMethods::Yield: { @@ -235,6 +350,10 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca } default: LOG_ERROR(HW_GPU, "Special puller engine method {:X} not implemented", method); + if (static bool dumped = false; g_nx2_ampere && !dumped) { + dumped = true; + dma_pusher.DumpCurrentEntry("unknown host method"); + } break; } } diff --git a/src/video_core/engines/puller.h b/src/video_core/engines/puller.h index e8994f5640..2b9e2c5df3 100644 --- a/src/video_core/engines/puller.h +++ b/src/video_core/engines/puller.h @@ -7,7 +7,9 @@ #pragma once #include +#include #include +#include #include #include "common/bit_field.h" #include "common/common_funcs.h" @@ -70,6 +72,24 @@ public: BitField<8, 24, u32> syncpoint_id; }; + /// Switch 2: a semaphore acquire this channel is suspended on. + struct PendingAcquire { + GPUVAddr address; + u32 value; + u32 operation; + std::chrono::steady_clock::time_point since; + bool reported; + + [[nodiscard]] bool Matches(u32 current) const; + }; + + [[nodiscard]] bool HasPendingAcquire() const { + return pending_acquire.has_value(); + } + [[nodiscard]] bool CanResumeAcquire(DmaPusher& dma_pusher); + /// Clears the pending acquire once it is met; false while the channel has to keep waiting. + bool ResumeAcquire(DmaPusher& dma_pusher); + void CallMethod(DmaPusher& dma_pusher, const MethodCall& method_call); void CallMultiMethod(DmaPusher& dma_pusher, u32 method, u32 subchannel, const u32* base_start, u32 amount, u32 methods_pending); void BindRasterizer(DmaPusher& dma_pusher, VideoCore::RasterizerInterface* rasterizer); @@ -125,8 +145,11 @@ private: void ProcessSemaphoreAcquire(DmaPusher& dma_pusher); void ProcessSemaphoreRelease(DmaPusher& dma_pusher); void ProcessSemaphoreTriggerMethod(DmaPusher& dma_pusher); + void ProcessSemaphoreExecute(DmaPusher& dma_pusher, u32 argument); [[nodiscard]] bool ExecuteMethodOnEngine(DmaPusher& dma_pusher, u32 method); + std::optional pending_acquire; + /// Mapping of command subchannels to their bound engine ids std::array bound_engines{}; diff --git a/src/video_core/gpu_thread.cpp b/src/video_core/gpu_thread.cpp index a01375e7be..97e18df499 100644 --- a/src/video_core/gpu_thread.cpp +++ b/src/video_core/gpu_thread.cpp @@ -4,6 +4,9 @@ // SPDX-FileCopyrightText: Copyright 2019 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include +#include + #include "common/assert.h" #include "common/scope_exit.h" #include "common/settings.h" @@ -36,7 +39,16 @@ void ThreadManager::StartThread(VideoCore::RendererBase& renderer, Core::Fronten auto current_context = context.Acquire(); CommandDataContainer next; while (!stop_token.stop_requested()) { - state.queue.PopWait(next, stop_token); + if (scheduler.HasPendingChannels()) { + // Suspended channels wait on semaphores the CPU releases: poll them meanwhile. + scheduler.ResumePending(system.GPU()); + if (!state.queue.TryPop(next)) { + std::this_thread::sleep_for(std::chrono::milliseconds{1}); + continue; + } + } else { + state.queue.PopWait(next, stop_token); + } if (stop_token.stop_requested()) { break; } diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 0290500e7c..52e9db189d 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -251,7 +251,10 @@ T MemoryManager::Read(GPUVAddr addr) const { return value; } - ASSERT(false); + static std::atomic reported{}; + if (reported.fetch_add(1) < 16) { + LOG_ERROR(HW_GPU, "Read{} of unmapped GPU address {:#x}", sizeof(T) * 8, addr); + } return {}; } diff --git a/src/video_core/nx2.h b/src/video_core/nx2.h index 9fdf0a7c68..ef6619df3f 100644 --- a/src/video_core/nx2.h +++ b/src/video_core/nx2.h @@ -3,6 +3,7 @@ #pragma once +#include #include #include "common/common_types.h" @@ -13,10 +14,34 @@ namespace Tegra { // Ampere classes and the Turing+ macro ISA, which Eden's Maxwell engines don't implement. inline std::atomic g_nx2_ampere{false}; -// NVN2 queues release their progress values into one GPU page and wait (GPU acquires and CPU -// polls) on the same offsets in the page below it. Something on the console copies one into -// the other; until that is known, the puller mirrors releases that land in the source page. -inline std::atomic g_nx2_semaphore_source_page{0}; // GPU VA of the released-to page -inline std::atomic g_nx2_semaphore_mirror_page{0}; // GPU VA of the waited-on page +// NVN2 queues release their progress values into the second page of a two-page GPU-accessible +// region and GPU acquires read the same offsets in the first page (the CPU polls the second). +// Something on the console copies one into the other; until that is known, the puller mirrors +// releases that land in a second page. Every queue has such a region, so keep a small table. +inline constexpr size_t Nx2MaxSemaphoreRegions = 32; +inline std::array, Nx2MaxSemaphoreRegions> g_nx2_semaphore_regions{}; // starts + +inline void Nx2AddSemaphoreRegion(u64 start) { + for (auto& region : g_nx2_semaphore_regions) { + u64 expected = 0; + if (region.load() == start || region.compare_exchange_strong(expected, start)) { + return; + } + } +} + +/// GPU VA in the first page that mirrors a release at `address`, or 0. +inline u64 Nx2SemaphoreMirror(u64 address) { + for (const auto& region : g_nx2_semaphore_regions) { + const u64 start = region.load(std::memory_order_relaxed); + if (start == 0) { + break; + } + if (address >= start + 0x1000 && address < start + 0x2000) { + return address - 0x1000; + } + } + return 0; +} } // namespace Tegra 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/engines/puller.cpp') 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/engines/puller.cpp') 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 64c91974f4455a63a025149a4ba8e2c46894bdc7 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 04:56:14 -0400 Subject: nx2: frame-locked input replay and quieter logs EDEN_NX2_KEY_RECORD records keyboard transitions as "f:code" against a global count of guest presents; EDEN_NX2_KEY_SCRIPT replays them on the same frames (or by seconds, as before), so a played session reproduces exactly. Debug aids: the PC/LR/SP of the first unmapped reads, NSO load addresses, EDEN_NX2_IPC_SUMMARY (request counts every 5 s), EDEN_NX2_FORCE_LOD0. Unmapped-access logging is rate limited and the per-draw/semaphore spam is demoted, so long sessions stay under the log size cap. --- src/core/arm/dynarmic/arm_dynarmic_64.cpp | 20 ++++++-- src/core/core.cpp | 2 + src/core/core.h | 4 ++ src/core/hle/service/nx2_gsv.h | 1 + src/core/hle/service/service.cpp | 57 ++++++++++++++++++++++ src/core/loader/nso.cpp | 1 + src/core/memory.cpp | 14 +++++- src/core/memory.h | 3 ++ src/video_core/engines/puller.cpp | 4 +- src/video_core/renderer_vulkan/vk_rasterizer.cpp | 2 +- .../renderer_vulkan/vk_texture_cache.cpp | 5 ++ src/yuzu_cmd/emu_window/emu_window_sdl3.cpp | 23 +++++++++ src/yuzu_cmd/yuzu.cpp | 26 ++++++---- 13 files changed, 145 insertions(+), 17 deletions(-) (limited to 'src/video_core/engines/puller.cpp') diff --git a/src/core/arm/dynarmic/arm_dynarmic_64.cpp b/src/core/arm/dynarmic/arm_dynarmic_64.cpp index 894857a86d..c2bb05afba 100644 --- a/src/core/arm/dynarmic/arm_dynarmic_64.cpp +++ b/src/core/arm/dynarmic/arm_dynarmic_64.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include "common/settings.h" #include "core/arm/dynarmic/arm_dynarmic.h" #include "core/arm/dynarmic/arm_dynarmic_64.h" @@ -24,13 +25,24 @@ DynarmicCallbacks64::DynarmicCallbacks64(ArmDynarmic64& parent, Kernel::KProcess u64 DynarmicCallbacks64::MemoryRead(u64 vaddr, size_t size) { CheckMemoryAccess(vaddr, size, Kernel::DebugWatchpointType::Read); + u64 value{}; switch (size) { - case sizeof(u64): return m_memory.Read64(vaddr); - case sizeof(u32): return m_memory.Read32(vaddr); - case sizeof(u16): return m_memory.Read16(vaddr); - case sizeof(u8): return m_memory.Read8(vaddr); + case sizeof(u64): value = m_memory.Read64(vaddr); break; + case sizeof(u32): value = m_memory.Read32(vaddr); break; + case sizeof(u16): value = m_memory.Read16(vaddr); break; + case sizeof(u8): value = m_memory.Read8(vaddr); break; default: UNREACHABLE(); } + // With memory aborts ignored the guest keeps running on zeros: say where it read from. + if (auto& unmapped = Core::Memory::UnmappedReadFlag(); unmapped != 0) { + unmapped = 0; + static std::atomic logged{}; + if (logged.fetch_add(1) < 32) { + LOG_ERROR(Core_ARM, "Unmapped read {:#x} ({} bytes) at pc {:#x} lr {:#x} sp {:#x}", vaddr, size, + m_parent.m_jit->GetPC(), m_parent.m_jit->GetRegister(30), m_parent.m_jit->GetSP()); + } + } + return value; } Dynarmic::A64::Vector DynarmicCallbacks64::MemoryRead128(u64 vaddr) { CheckMemoryAccess(vaddr, 16, Kernel::DebugWatchpointType::Read); diff --git a/src/core/core.cpp b/src/core/core.cpp index 06543b0833..eee91ef4fc 100644 --- a/src/core/core.cpp +++ b/src/core/core.cpp @@ -68,6 +68,8 @@ namespace Core { +std::atomic g_nx2_present_count{}; + FileSys::VirtualFile GetGameFileFromPath(const FileSys::VirtualFilesystem& vfs, const std::string& path) { // To account for split 00+01+etc files. diff --git a/src/core/core.h b/src/core/core.h index d88ed64151..0276679d76 100644 --- a/src/core/core.h +++ b/src/core/core.h @@ -6,6 +6,7 @@ #pragma once +#include #include #include #include @@ -115,6 +116,9 @@ class RenderdocAPI; namespace Core { +/// Frames the guest has presented (Switch 2 vi2 present calls); input replay is locked to it. +extern std::atomic g_nx2_present_count; + class CpuManager; class Debugger; class DeviceMemory; diff --git a/src/core/hle/service/nx2_gsv.h b/src/core/hle/service/nx2_gsv.h index 89f3a9fd50..302c5a41a8 100644 --- a/src/core/hle/service/nx2_gsv.h +++ b/src/core/hle/service/nx2_gsv.h @@ -564,6 +564,7 @@ private: void PresentClientBuffer(HLERequestContext& ctx) { const auto* w = ctx.CommandBuffer(); ++m_presents; + Core::g_nx2_present_count.fetch_add(1, std::memory_order_relaxed); if (m_presents <= 4 || (m_presents % 600) == 0) { LOG_INFO(Service, "nx2:vi2 229 present #{} token={:#x} slot={} image={}", m_presents, u64{w[2]} | (u64{w[3]} << 32), w[4], w[5]); diff --git a/src/core/hle/service/service.cpp b/src/core/hle/service/service.cpp index 931de1c3ef..df929e8193 100644 --- a/src/core/hle/service/service.cpp +++ b/src/core/hle/service/service.cpp @@ -4,6 +4,13 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include "common/scope_exit.h" +#include +#include +#include +#include +#include +#include #include #include #include "common/assert.h" @@ -125,9 +132,59 @@ void ServiceFrameworkBase::InvokeRequestTipc(HLERequestContext& ctx) { handler_invoker(this, info->handler_callback, ctx); } +namespace { +// EDEN_NX2_IPC_SUMMARY: every 5 s, log how often each service command ran and with which result +// since the last summary (finds a game polling something that never becomes ready). +void CountRequest(const std::string& service, u32 command, u32 type, u32 result) { + static const bool enabled = std::getenv("EDEN_NX2_IPC_SUMMARY") != nullptr; + if (!enabled) { + return; + } + static std::mutex mutex; + static std::map, u32> counts; + static auto last = std::chrono::steady_clock::now(); + std::scoped_lock lock{mutex}; + ++counts[{service, command, type, result}]; + const auto now = std::chrono::steady_clock::now(); + if (now - last < std::chrono::seconds{5}) { + return; + } + last = now; + std::vector> lines; + for (const auto& [key, count] : counts) { + lines.emplace_back(count, fmt::format("{}::{}{} -> {:#x}", std::get<0>(key), std::get<1>(key), + std::get<2>(key) >= 16 ? " (tipc)" : "", + std::get<3>(key))); + } + std::sort(lines.begin(), lines.end(), [](const auto& a, const auto& b) { return a.first > b.first; }); + std::string text; + for (size_t i = 0; i < std::min(lines.size(), 24); ++i) { + text += fmt::format(" | {}x {}", lines[i].first, lines[i].second); + } + LOG_INFO(Service, "ipc summary ({} kinds){}", lines.size(), text); + counts.clear(); +} +} // Anonymous namespace + Result ServiceFrameworkBase::HandleSyncRequest(Kernel::KServerSession& session, HLERequestContext& ctx) { const auto guard = LockService(); + SCOPE_EXIT { + const auto* cmd_buf = ctx.CommandBuffer(); + u32 result_word = 0; + const u32 type = static_cast(ctx.GetCommandType()); + if (type >= 16) { + result_word = cmd_buf[2]; + } else { + for (int i = 0; i < 12; ++i) { + if (cmd_buf[i] == 0x4F434653) { // "SFCO" + result_word = cmd_buf[i + 2]; + break; + } + } + } + CountRequest(service_name, ctx.GetCommand(), type, result_word); + }; Result result = ResultSuccess; diff --git a/src/core/loader/nso.cpp b/src/core/loader/nso.cpp index ea1d1c9ecd..7e3b9a54c2 100644 --- a/src/core/loader/nso.cpp +++ b/src/core/loader/nso.cpp @@ -247,6 +247,7 @@ std::optional AppLoader_NSO::LoadModule(Kernel::KProcess& process, Core:: } // Load codeset for current process + LOG_INFO(Loader, "Module {} loaded at {:#x} (size {:#x})", nso_file.GetName(), load_base, image_size); process.LoadModule(system.Kernel(), std::move(codeset), load_base); return load_base + image_size; } diff --git a/src/core/memory.cpp b/src/core/memory.cpp index eabc6dfccb..6e0ab0ee6d 100644 --- a/src/core/memory.cpp +++ b/src/core/memory.cpp @@ -5,6 +5,7 @@ // SPDX-FileCopyrightText: 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include #include #include #include @@ -36,6 +37,11 @@ namespace Core::Memory { +u32& UnmappedReadFlag() { + thread_local u32 flag{}; + return flag; +} + static inline bool AddressSpaceContains(const Common::PageTable& table, const Common::ProcessAddress addr, const std::size_t size) { const Common::ProcessAddress max_addr = 1ULL << table.GetAddressSpaceBits(); return addr + size >= addr && addr + size <= max_addr; @@ -623,7 +629,13 @@ struct Memory::Impl { inline T Read(Common::ProcessAddress vaddr) noexcept requires(std::is_trivially_copyable_v) { const u64 addr = GetInteger(vaddr); if (auto const ptr = GetPointerImpl(addr, [addr]() { - LOG_ERROR(HW_Memory, "Unmapped Read{} @ {:#016x}", sizeof(T) * 8, addr); + // A guest spinning on a bad pointer reads unmapped memory tens of thousands of times + // a second: keep the first ones and then a sample. + static std::atomic unmapped_reads{}; + if (const u64 count = unmapped_reads.fetch_add(1); count < 1000 || count % 100000 == 0) { + LOG_ERROR(HW_Memory, "Unmapped Read{} @ {:#016x} (#{})", sizeof(T) * 8, addr, count); + } + ++UnmappedReadFlag(); }, [&]() { HandleRasterizerDownload(addr, sizeof(T)); }); ptr) [[likely]] { diff --git a/src/core/memory.h b/src/core/memory.h index 7167efbb84..204b7514f6 100644 --- a/src/core/memory.h +++ b/src/core/memory.h @@ -38,6 +38,9 @@ class MemoryManager; namespace Core::Memory { +/// Counts unmapped reads on this host thread since the CPU callbacks last checked (debug aid). +u32& UnmappedReadFlag(); + /** * Page size used by the ARM architecture. This is the smallest granularity with which memory can * be mapped. diff --git a/src/video_core/engines/puller.cpp b/src/video_core/engines/puller.cpp index f4aaed6036..54abf790d2 100644 --- a/src/video_core/engines/puller.cpp +++ b/src/video_core/engines/puller.cpp @@ -335,7 +335,7 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca break; } case BufferMethods::MemOpA: { - LOG_ERROR(HW_GPU, "Memory Operation A"); + LOG_TRACE(HW_GPU, "Memory Operation A"); break; } case BufferMethods::MemOpB: { @@ -345,7 +345,7 @@ void Puller::CallPullerMethod(DmaPusher& dma_pusher, const MethodCall& method_ca } case BufferMethods::MemOpC: case BufferMethods::MemOpD: { - LOG_ERROR(HW_GPU, "Memory Operation C,D"); + LOG_TRACE(HW_GPU, "Memory Operation C,D"); break; } case BufferMethods::SemaphoreAcquire: { diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index 1b05716653..1b0493c4b5 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -250,7 +250,7 @@ void RasterizerVulkan::PrepareDraw(bool is_indexed, Func&& draw_func) { auto& native = maxwell3d->GetNativeMacroEngine(); if (native.native_turing && !submitted) { native.native_execution_failed = true; - LOG_ERROR(HW_GPU, "Native graphics draw failed; refusing submission completion"); + LOG_DEBUG(HW_GPU, "Native graphics draw failed; refusing submission completion"); } }; FlushWork(); diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.cpp b/src/video_core/renderer_vulkan/vk_texture_cache.cpp index 29afaac01b..312617475d 100644 --- a/src/video_core/renderer_vulkan/vk_texture_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_texture_cache.cpp @@ -2826,6 +2826,11 @@ Sampler::Sampler(TextureCacheRuntime& runtime, const Tegra::Texture::TSCEntry& t if (tsc.mipmap_filter != TextureMipmapFilter::None) { min_lod = tsc.MinLod(); max_lod = tsc.MaxLod(); + // EDEN_NX2_FORCE_LOD0: sample only the top mip level (debug aid for mip layout problems). + if (static const bool force_lod0 = std::getenv("EDEN_NX2_FORCE_LOD0") != nullptr; force_lod0) { + min_lod = 0.0f; + max_lod = 0.25f; + } } base_ci = VkSamplerCreateInfo{ .sType = VK_STRUCTURE_TYPE_SAMPLER_CREATE_INFO, diff --git a/src/yuzu_cmd/emu_window/emu_window_sdl3.cpp b/src/yuzu_cmd/emu_window/emu_window_sdl3.cpp index f2026e9722..d62acb36cd 100644 --- a/src/yuzu_cmd/emu_window/emu_window_sdl3.cpp +++ b/src/yuzu_cmd/emu_window/emu_window_sdl3.cpp @@ -4,6 +4,10 @@ // SPDX-FileCopyrightText: 2016 Citra Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include +#include +#include +#include #include #include #ifdef __EMSCRIPTEN__ @@ -118,7 +122,26 @@ void EmuWindow_SDL3::OnFingerUp() { input_subsystem->GetTouchScreen()->ReleaseAllTouch(); } +// Seconds since the first frame; shared by the key script (yuzu.cpp) and key recording. +double Nx2KeyClockSeconds() { + static const auto begin = std::chrono::steady_clock::now(); + return std::chrono::duration(std::chrono::steady_clock::now() - begin).count(); +} + void EmuWindow_SDL3::OnKeyEvent(int key, u8 state) { + // EDEN_NX2_KEY_RECORD=: append each press/release ("seconds:code", "seconds:-code"). + static std::FILE* const record = [] { + const char* path = std::getenv("EDEN_NX2_KEY_RECORD"); + return path ? std::fopen(path, "w") : nullptr; + }(); + static std::array held{}; + if (record && key >= 0 && key < static_cast(held.size()) && held[key] != (state != 0)) { + held[key] = state != 0; + std::fprintf(record, "f%llu:%s%d,", + static_cast(Core::g_nx2_present_count.load()), + state != 0 ? "" : "-", key); + std::fflush(record); + } if (state != 0) { input_subsystem->GetKeyboard()->PressKey(static_cast(key)); } else { diff --git a/src/yuzu_cmd/yuzu.cpp b/src/yuzu_cmd/yuzu.cpp index b46534ec9d..96146abec3 100644 --- a/src/yuzu_cmd/yuzu.cpp +++ b/src/yuzu_cmd/yuzu.cpp @@ -333,9 +333,12 @@ extern "C" SDL_AppResult SDL_AppInit(void **appstate, int argc, char **argv) { } // 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. +double Nx2KeyClockSeconds(); // emu_window_sdl3.cpp + static void RunKeyScript(InputCommon::InputSubsystem& input_subsystem) { struct KeyEvent { - double seconds; + double seconds; // when frame is 0 + u64 frame; // "f:code": after that many presents int code; }; static const std::vector script = [] { @@ -347,22 +350,27 @@ static void RunKeyScript(InputCommon::InputSubsystem& input_subsystem) { 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))}); + const std::string when = item.substr(0, colon); + const int code = std::stoi(item.substr(colon + 1)); + if (!when.empty() && when[0] == 'f') { + events.push_back({0.0, std::stoull(when.substr(1)), code}); + } else { + events.push_back({std::stod(when), 0, code}); + } } 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 double elapsed = Nx2KeyClockSeconds(); + const u64 frame = Core::g_nx2_present_count.load(); + while (next < script.size() && (script[next].frame != 0 ? script[next].frame <= frame + : 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); + LOG_INFO(Frontend, "Key script: {} key {} at {:.1f} s (frame {})", + code < 0 ? "release" : "press", code < 0 ? -code : code, elapsed, frame); if (code < 0) { input_subsystem.GetKeyboard()->ReleaseKey(static_cast(-code)); } else { -- 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/engines/puller.cpp') 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