From 51207c73c889dbdc4a586a5da23314f6d33e4eb9 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Tue, 27 Oct 2026 22:26:16 -0400 Subject: Switch 2: Turing macros, native compute, Ampere shader lowering, native tiling From changes.patch: the TU104/GA10x MME interpreter (shared macro context for the process's channels, also on the compute class), native compute with QMD 3.0 dependency scheduling, the Ampere-to-Maxwell shader lowering, native vertex stream sizes, and the console's GOB swizzle for textures, DMA copies and the host unswizzle shaders. Switch 2 channels turn all of it on at init. The lowering is widened for title 4 (all 32 of its programs lower): general TEX (any dimension/LOD/mask, operand split 1+1/2+1/2+2/4+n, any but consistent descriptor buffer), cube/array, SULD channel masks, 128-bit STG, indexed IPA, VOTE, SHFL with any immediate or predicate output, R2P/P2R byte select, uniform ULOP3/ULEA/UBREV/UFLO.SH/UISETP signed/UMOV URZ, IMNMX cbuf, IADD3 and FFMA uniform operands, FMUL /4 /8 x8, and BSSY/BSYNC/BREAK/WARPSYNC as no-ops. The capture stops at unmapped memory. EDEN_NX2_SHADER_DUMP dumps every native program; tools in C:\dev\nx2work test the lowering offline. --- src/video_core/shader_cache.cpp | 30 +++++++++++++++++++++++------- 1 file changed, 23 insertions(+), 7 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index f1aff2810f..75e5432ab2 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -11,6 +11,7 @@ #include "common/assert.h" #include "shader_recompiler/frontend/maxwell/control_flow.h" +#include "shader_recompiler/exception.h" #include "shader_recompiler/object_pool.h" #include "video_core/control/channel_state.h" #include "video_core/dirty_flags.h" @@ -43,7 +44,7 @@ void ShaderCache::SyncGuestHost() { ShaderCache::ShaderCache(Tegra::MaxwellDeviceMemoryManager& device_memory_) : device_memory{device_memory_} {} -bool ShaderCache::RefreshStages(std::array& unique_hashes) { +bool ShaderCache::RefreshStages(std::array& unique_hashes) try { auto& dirty{maxwell3d->dirty.flags}; if (!dirty[VideoCommon::Dirty::Shaders]) { return last_shaders_valid; @@ -51,6 +52,7 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { dirty[VideoCommon::Dirty::Shaders] = false; const GPUVAddr base_addr{maxwell3d->regs.program_region.Address()}; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; for (size_t index = 0; index < Tegra::Engines::Maxwell3D::Regs::MaxShaderProgram; ++index) { if (!maxwell3d->regs.IsShaderConfigEnabled(index)) { unique_hashes[index] = 0; @@ -63,7 +65,9 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { unique_hashes[index] = 0; continue; } - const GPUVAddr shader_addr{base_addr + shader_config.offset}; + const GPUVAddr shader_addr{native + ? (static_cast(shader_config.reserved[0] & 0xff) << 32) | shader_config.reserved[1] + : base_addr + shader_config.offset}; const std::optional cpu_shader_addr{gpu_memory->GpuToCpuAddress(shader_addr)}; if (!cpu_shader_addr) { LOG_ERROR(HW_GPU, "Invalid GPU address for shader {:#016x}", shader_addr); @@ -72,8 +76,8 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { } const ShaderInfo* shader_info{TryGet(*cpu_shader_addr)}; if (!shader_info) { - const u32 start_address{shader_config.offset}; - GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, base_addr, start_address}; + const u32 start_address{native ? 0 : shader_config.offset}; + GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, native ? shader_addr : base_addr, start_address}; shader_info = MakeShaderInfo(env, *cpu_shader_addr); } shader_infos[index] = shader_info; @@ -81,6 +85,13 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) { } last_shaders_valid = true; return true; +} catch (const Shader::Exception& error) { + auto& native = maxwell3d->GetNativeMacroEngine(); + if (!native.native_turing) throw; + native.native_execution_failed = true; + last_shaders_valid = false; + LOG_ERROR(HW_GPU, "Native graphics shader rejected: {}", error.what()); + return false; } const ShaderInfo* ShaderCache::ComputeShader() { @@ -103,14 +114,19 @@ void ShaderCache::GetGraphicsEnvironments(GraphicsEnvironments& result, const std::array& unique_hashes) { size_t env_index{}; const GPUVAddr base_addr{maxwell3d->regs.program_region.Address()}; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; for (size_t index = 0; index < NUM_PROGRAMS; ++index) { if (unique_hashes[index] == 0) { continue; } const auto program{static_cast(index)}; auto& env{result.envs[index]}; - const u32 start_address{maxwell3d->regs.pipelines[index].offset}; - env = GraphicsEnvironment{*maxwell3d, *gpu_memory, program, base_addr, start_address}; + const auto& config = maxwell3d->regs.pipelines[index]; + const u32 start_address{native ? 0 : config.offset}; + const GPUVAddr address = native + ? (static_cast(config.reserved[0] & 0xff) << 32) | config.reserved[1] + : base_addr; + env = GraphicsEnvironment{*maxwell3d, *gpu_memory, program, address, start_address}; env.SetCachedSize(shader_infos[index]->size_bytes); result.env_ptrs[env_index++] = &env; } @@ -241,7 +257,7 @@ const ShaderInfo* ShaderCache::MakeShaderInfo(GenericEnvironment& env, VAddr cpu info->unique_hash = env.CalculateHash(); info->size_bytes = env.ReadSizeBytes(); } - const size_t size_bytes{info->size_bytes}; + const size_t size_bytes{env.SourceSizeBytes() ? env.SourceSizeBytes() : info->size_bytes}; const ShaderInfo* const result{info.get()}; Register(std::move(info), cpu_addr, size_bytes); return result; -- cgit v1.2.3 From 869c6e4b33a6432166597f6f48812d44b3f5de02 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 01:18:23 -0400 Subject: nx2: lower every shader title 4 uses in game - Tensor-core ops for its neural compute pass: HMMA.16816.F16 (SHFL + HFMA2), MOVM.16.MT88 and LDSM .MT88. Opt-in via EDEN_NX2_TENSOR: the expansion (372 HMMAs) takes the host driver minutes to compile. - HSETP2 with an immediate half pair; skip @!PT instructions. - Constant bank 24, read only by three NVN2 fullscreen vertex shaders and never bound by the game, reads from bank 17 (also never bound: zeros, as on hw). - Debug dumps of the vertex stage's bound constant buffers on rejection. All 214 shaders captured after the menu lower (213 without EDEN_NX2_TENSOR). --- .../frontend/ampere/lower_to_maxwell.h | 144 +++++++++++++++++++-- src/video_core/shader_cache.cpp | 17 +++ 2 files changed, 152 insertions(+), 9 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index fc3f088c57..d196a1d1fc 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -3,6 +3,7 @@ // Encoding reference: Mesa NAK sm70_encode.rs (MIT), Copyright 2022 Collabora, Ltd. #pragma once +#include #include #include #include @@ -79,6 +80,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { auto insn = code[index]; result.source_instruction_offsets.push_back(static_cast(result.instructions.size())); u32 opcode = insn.Bits(0, 12); + if (insn.Bits(12, 4) == 15) continue; // @!PT: never executes (compiler padding) std::optional active_link; u32 active_entry = 0; if (main_exited) { @@ -130,11 +132,17 @@ inline LoweredProgram LowerToMaxwell(std::span code) { u64 predicate = u64{guarded_uniform ? ((native_predicate & 8) | 6) : native_predicate} << 16; const u64 dest = insn.Bits(16, 8); const auto base = [&](u64 opcode_bits) { return opcode_bits | predicate | dest; }; + // NVN2 fullscreen vertex shaders read bank 24, which the game never binds (the console reads + // zeros from an unbound bank); Maxwell has 18, so read it from bank 17, also never bound. + const auto bank = [&]() -> u32 { + const u32 value = insn.Bits(54, 5); + return value == 24 ? 17U : value; + }; const auto cbuf = [&]() -> u64 { - return (u64{insn.Bits(38, 16) / 4} << 20) | (u64{insn.Bits(54, 5)} << 34); + return (u64{insn.Bits(38, 16) / 4} << 20) | (u64{bank()} << 34); }; const auto valid_cbuf = [&] { - return !insn.Bits(91, 1) && insn.Bits(54, 5) < 18 && + return !insn.Bits(91, 1) && bank() < 18 && (insn.Bits(38, 16) % 4) == 0; }; std::string_view operand_error; @@ -575,7 +583,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { const u32 width = insn.Bits(73, 3); if (uniform_dest >= 16 || (width != 4 && width != 5) || (width == 5 && ((uniform_dest & 1) || uniform_dest >= 15)) || - insn.Bits(54, 5) >= 18 || insn.Bits(38, 16) % 4 || insn.Bits(62, 2) || + bank() >= 18 || insn.Bits(38, 16) % 4 || insn.Bits(62, 2) || (insn.words[2] & ~0xe00U) != (opcode == 0xab9 ? 0U : 0x08000000U) || (opcode == 0xabb && index_reg >= 16)) return fail(index, "native uniform constant load mode"); @@ -584,7 +592,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { } else { emit(0xef90000000000000ULL | predicate | (uniform_dest + 216) | (u64{opcode == 0xab9 ? 255U : index_reg + 216} << 8) | (u64{insn.Bits(38, 16)} << 20) | - (u64{insn.Bits(54, 5)} << 36) | (u64{width} << 48)); + (u64{bank()} << 36) | (u64{width} << 48)); } break; } @@ -607,10 +615,10 @@ inline LoweredProgram LowerToMaxwell(std::span code) { case 0xb82: { // LDC.64 indexed byte offset const u32 width = insn.Bits(73, 3); if ((width != 4 && width != 5) || insn.Bits(78, 2) || insn.Bits(91, 1) || - insn.Bits(54, 5) >= 18 || + bank() >= 18 || (width == 5 && ((dest & 1) || dest >= 254))) return fail(index, "native indexed constant load mode"); emit(base(0xef90000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | - (u64{insn.Bits(38, 16)} << 20) | (u64{insn.Bits(54, 5)} << 36) | (u64{width} << 48)); + (u64{insn.Bits(38, 16)} << 20) | (u64{bank()} << 36) | (u64{width} << 48)); break; } case 0xf89: { // SHFL with immediate lane and segmentation/clamp mask, any mode. @@ -852,18 +860,125 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{width} << 48)); break; } + case 0x23a: { // MOVM.16.MT88: transpose an 8x8 16-bit matrix held one row pair per lane. + // Lane t holds row t/4, columns 2(t%4)..+1; the result is the transposed matrix in the + // same layout, so each half comes from lane 8(t%4) + t/8 (+4) at half (t/4)%2. + const u32 source = insn.Bits(24, 8); + if (insn.words[1] || insn.words[2] || source >= 246 || dest >= 246) + return fail(index, "native matrix move mode"); + emit(0xf0c8000000000000ULL | predicate | 246); // laneid + emit(0x0400000000000000ULL | predicate | 249 | (u64{246} << 8) | (u64{3} << 20)); + emit(0x3848000000000000ULL | predicate | 249 | (u64{249} << 8) | (u64{3} << 20)); // 8(t%4) + emit(0x3828000000000000ULL | predicate | 247 | (u64{246} << 8) | (u64{3} << 20)); // t/8 + emit(0x5c10000000000000ULL | predicate | 249 | (u64{249} << 8) | (u64{247} << 20)); + emit(0xef10000000000000ULL | predicate | 247 | (u64{source} << 8) | (u64{249} << 20) | + (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x1c00000000000000ULL | predicate | 249 | (u64{249} << 8) | (u64{4} << 20)); + emit(0xef10000000000000ULL | predicate | 248 | (u64{source} << 8) | (u64{249} << 20) | + (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x3828000000000000ULL | predicate | 249 | (u64{246} << 8) | (u64{2} << 20)); + emit(0x0400000000000000ULL | predicate | 249 | (u64{249} << 8) | (u64{1} << 20)); + emit(0x3848000000000000ULL | predicate | 249 | (u64{249} << 8) | (u64{4} << 20)); // 0 or 16 + emit(0x5c28000000000000ULL | predicate | 247 | (u64{247} << 8) | (u64{249} << 20)); + emit(0x0400000000000000ULL | predicate | 247 | (u64{247} << 8) | (u64{0xffff} << 20)); + emit(0x5c28000000000000ULL | predicate | 248 | (u64{248} << 8) | (u64{249} << 20)); + emit(0x5c18000000000000ULL | predicate | dest | (u64{248} << 8) | (u64{247} << 20) | (u64{16} << 39)); + break; + } + case 0x23c: { // HMMA.16816.F16: D(16x8) = A(16x16) * B(16x8) + C, f16 fragments. + // Lane t (group g = t/4, i = t%4): A regs hold rows g / g+8 at k 2i..2i+1 (a0, a1) and + // 8+2i.. (a2, a3); B regs hold k 2i.. (b0) / 8+2i.. (b1) of column g; D holds rows + // g / g+8 at columns 2i, 2i+1. Rebuilt from shuffles with HFMA2, accumulating in R252/R253. + // The expansion is correct but large: title 4's 372-HMMA network takes the host driver + // minutes to compile, so it stays opt-in (the QMD still completes when rejected). + static const bool enabled = std::getenv("EDEN_NX2_TENSOR") != nullptr; + if (!enabled) return fail(index, "native matrix multiply (set EDEN_NX2_TENSOR to emulate)"); + const u32 a = insn.Bits(24, 8), b = insn.Bits(32, 8), c = insn.Bits(64, 8); + if ((insn.words[1] & ~0xffU) || (insn.words[2] & ~0xffU) != 0x800 || dest + 2 > 246 || + a + 4 > 246 || b + 2 > 246 || (c != 255 && c + 2 > 246)) + return fail(index, "native matrix multiply mode"); + const auto shuffle = [&](u32 target, u32 source_reg) { + emit(0xef10000000000000ULL | predicate | target | (u64{source_reg} << 8) | (u64{246} << 20) | + (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + }; + const auto lane_a = [&](u32 i) { // 4g + i + emit(0xf0c8000000000000ULL | predicate | 246); + emit(0x0400000000000000ULL | predicate | 246 | (u64{246} << 8) | (u64{0x1c} << 20)); + if (i) emit(0x1c00000000000000ULL | predicate | 246 | (u64{246} << 8) | (u64{i} << 20)); + }; + const auto lane_b = [&](u32 add) { // 8(t%4) + add: column 2(t%4) (+1 when add >= 4) + emit(0xf0c8000000000000ULL | predicate | 246); + emit(0x0400000000000000ULL | predicate | 246 | (u64{246} << 8) | (u64{3} << 20)); + emit(0x3848000000000000ULL | predicate | 246 | (u64{246} << 8) | (u64{3} << 20)); + if (add) emit(0x1c00000000000000ULL | predicate | 246 | (u64{246} << 8) | (u64{add} << 20)); + }; + const auto hfma2 = [&](u32 accumulator, u32 swizzle_a, u32 pair) { + emit(0x5d00000000000000ULL | predicate | accumulator | (u64{251} << 8) | (u64{pair} << 20) | + (u64{accumulator} << 39) | (u64{swizzle_a} << 47)); + }; + for (u32 row = 0; row < 2; ++row) { + if (c == 255) mov_imm(252 + row, 0); + else emit(0x5c98000000000000ULL | predicate | (252 + row) | (u64{15} << 39) | (u64{c + row} << 20)); + } + for (u32 i = 0; i < 4; ++i) { + for (u32 k_block = 0; k_block < 2; ++k_block) { + lane_b(i); + shuffle(247, b + k_block); + lane_b(4 + i); + shuffle(248, b + k_block); + // P0 = (B[2i][2t], B[2i][2t+1]), P1 = (B[2i+1][2t], B[2i+1][2t+1]) + emit(0x0400000000000000ULL | predicate | 249 | (u64{247} << 8) | (u64{0xffff} << 20)); + emit(0x5c18000000000000ULL | predicate | 249 | (u64{248} << 8) | (u64{249} << 20) | (u64{16} << 39)); + emit(0x3828000000000000ULL | predicate | 250 | (u64{247} << 8) | (u64{16} << 20)); + emit(0x0400000000000000ULL | predicate | 248 | (u64{248} << 8) | (u64{0xffff0000} << 20)); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{248} << 20)); + for (u32 row = 0; row < 2; ++row) { + lane_a(i); + shuffle(251, a + k_block * 2 + row); + hfma2(252 + row, 2, 249); // A.H0 broadcast * P0 + hfma2(252 + row, 3, 250); // A.H1 broadcast * P1 + } + } + } + for (u32 row = 0; row < 2; ++row) + emit(0x5c98000000000000ULL | predicate | (dest + row) | (u64{15} << 39) | (u64{252 + row} << 20)); + 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) || + const bool transpose = insn.Bits(78, 1) != 0; + if (insn.Bits(72, 2) == 3 || (insn.words[2] & ~0x4300U) || 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(0x5c10000000000000ULL | predicate | 253 | (u64{address} << 8) | (u64{255} << 20)); + address = 253; + } + if (transpose) { + // .MT88: lane t receives stored rows 2(t%4) and 2(t%4)+1 at column t/4 (16 bits each). + emit(0xf0c8000000000000ULL | predicate | 246); // laneid + emit(0x0400000000000000ULL | predicate | 247 | (u64{246} << 8) | (u64{3} << 20)); + emit(0x3848000000000000ULL | predicate | 247 | (u64{247} << 8) | (u64{1} << 20)); // row 2(t%4) + emit(0x3828000000000000ULL | predicate | 248 | (u64{246} << 8) | (u64{2} << 20)); + emit(0x3848000000000000ULL | predicate | 248 | (u64{248} << 8) | (u64{1} << 20)); // byte 2(t/4) + for (u32 matrix = 0; matrix < count; ++matrix) { + for (u32 half = 0; half < 2; ++half) { + emit(0x1c00000000000000ULL | predicate | 249 | (u64{247} << 8) | + (u64{matrix * 8 + half} << 20)); + emit(0xef10000000000000ULL | predicate | 250 | (u64{address} << 8) | (u64{249} << 20) | + (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{248} << 20)); + emit(0xef48000000000000ULL | predicate | (251 + half) | (u64{250} << 8) | + (offset << 20) | (u64{2} << 48)); // LDS.U16 + } + emit(0x5c18000000000000ULL | predicate | (dest + matrix) | (u64{252} << 8) | + (u64{251} << 20) | (u64{16} << 39)); + } + break; } emit(0xf0c8000000000000ULL | predicate | 250); // S2R laneid emit(0x3828000000000000ULL | predicate | 251 | (u64{250} << 8) | (u64{2} << 20)); // row @@ -1464,6 +1579,17 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{insn.Bits(80, 1)} << 55)); break; } + case 0x434: { // HSETP2 with an immediate half pair: A swizzle 74-75 (neg 72, abs 73), + // combine 69-70, H_AND 71, compare 76-79, FTZ 80, P/Q 81-86, input 87-90. + mov_imm(250, insn.words[1]); + emit(0x5d20000000000000ULL | predicate | insn.Bits(84, 3) | (u64{insn.Bits(81, 3)} << 3) | + (u64{insn.Bits(80, 1)} << 6) | (u64{insn.Bits(24, 8)} << 8) | (u64{250} << 20) | + (u64{insn.Bits(76, 4)} << 35) | (u64{insn.Bits(87, 3)} << 39) | + (u64{insn.Bits(90, 1)} << 42) | (u64{insn.Bits(72, 1)} << 43) | + (u64{insn.Bits(73, 1)} << 44) | (u64{insn.Bits(69, 2)} << 45) | + (u64{insn.Bits(74, 2)} << 47) | (u64{insn.Bits(71, 1)} << 49)); + break; + } case 0x20b: case 0x80b: case 0xa0b: // FSETP case 0x20c: case 0x80c: case 0xa0c: case 0xc0c: { // ISETP const bool floating = (opcode & 0x1ff) == 0xb; diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index 75e5432ab2..1f826313a1 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -5,6 +5,7 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include #include #include #include @@ -91,6 +92,22 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) try { native.native_execution_failed = true; last_shaders_valid = false; LOG_ERROR(HW_GPU, "Native graphics shader rejected: {}", error.what()); + // Once: show what the vertex stage has bound, to find where an unbound bank's data lives. + static std::vector dumped_errors; + if (const std::string what{error.what()}; + what.find("constant buffer") != std::string::npos && dumped_errors.size() < 12 && + std::count(dumped_errors.begin(), dumped_errors.end(), what) < 3) { + dumped_errors.push_back(what); + LOG_INFO(HW_GPU, "VS cbuf dump for: {}", what.substr(0, 40)); + for (u32 slot = 3; slot < 8; ++slot) { + const auto& buffer = maxwell3d->state.shader_stages[0].const_buffers[slot]; + if (!buffer.enabled) continue; + std::array words{}; + gpu_memory->ReadBlockUnsafe(buffer.address, words.data(), sizeof(words)); + LOG_INFO(HW_GPU, "VS cbuf {} at {:#x} size {:#x}: {:08x}", slot, buffer.address, buffer.size, + fmt::join(words, " ")); + } + } return false; } -- cgit v1.2.3 From e7cc2cea290280e519ce011049063a9e40e7cc9c Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 14:31:00 -0400 Subject: Cache rejected native shaders and log each rejection once A rejected program was lowered again (and logged) on every draw or dispatch: 1.7M log lines in 13 minutes of title 1. The shader cache keeps the first exception per address until guest memory there is written, and the error logs are deduplicated. --- src/video_core/engines/kepler_compute.cpp | 17 +++++++++++++++-- src/video_core/shader_cache.cpp | 30 ++++++++++++++++++++++++++++-- src/video_core/shader_cache.h | 6 ++++++ 3 files changed, 49 insertions(+), 4 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 8a3855468f..39c033d989 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -10,6 +10,9 @@ #include #include #include +#include +#include +#include #include "common/assert.h" #include "common/logging.h" #include "shader_recompiler/exception.h" @@ -316,11 +319,18 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar return true; } native_macro->native_execution_failed = true; - LOG_ERROR(HW_GPU, "Native QMD failed: {}", error); std::array qmd{}; if (memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { memory_manager.ReadBlockUnsafe(qmd_address, qmd.data(), sizeof(qmd)); } + // Log a failing program's first few dispatches only (QMD words 48-49 hold its code + // address); it fails the same way every frame. + static std::unordered_map failures; + const u64 program = (u64{qmd[49]} << 32) | qmd[48]; + if (failures[program ^ std::hash{}(std::string{error})]++ >= 2) { + return true; + } + LOG_ERROR(HW_GPU, "Native QMD failed: {}", error); if (qmd_diagnostics && memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { std::ofstream dump("native-qmd.bin", std::ios::binary); dump.write(reinterpret_cast(qmd.data()), sizeof(qmd)); @@ -438,7 +448,10 @@ bool KeplerCompute::DispatchNativeCompute(const NativeQmdScheduler::Descriptor& try { rasterizer->DispatchCompute(); } catch (const Shader::Exception& error) { - LOG_ERROR(HW_GPU, "Native compute shader rejected: {}", error.what()); + static std::unordered_set logged; + if (logged.insert(error.what()).second) { + LOG_ERROR(HW_GPU, "Native compute shader rejected: {}", error.what()); + } if (std::string_view{error.what()}.find("matrix multiply") != std::string_view::npos) { DescribeNativeResources(qmd); // The console runs this on tensor cores; stand in with a plain upscale of its input. diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index 1f826313a1..f92fe02068 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -5,6 +5,8 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include +#include #include #include #include @@ -91,7 +93,11 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) try { if (!native.native_turing) throw; native.native_execution_failed = true; last_shaders_valid = false; - LOG_ERROR(HW_GPU, "Native graphics shader rejected: {}", error.what()); + // Each rejected program is logged once (the draws that use it keep failing quietly). + static std::unordered_set logged; + if (logged.insert(error.what()).second) { + LOG_ERROR(HW_GPU, "Native graphics shader rejected: {}", error.what()); + } // Once: show what the vertex stage has bound, to find where an unbound bank's data lives. static std::vector dumped_errors; if (const std::string what{error.what()}; @@ -177,6 +183,12 @@ void ShaderCache::Register(std::unique_ptr data, VAddr addr, size_t void ShaderCache::InvalidatePagesInRegion(VAddr addr, size_t size) { const VAddr addr_end = addr + size; + { + std::scoped_lock lock{rejected_mutex}; + boost::unordered::erase_if(rejected_native, [&](const auto& rejected) { + return rejected.first >= addr && rejected.first < addr_end; + }); + } const u64 page_end = (addr_end + YUZU_PAGESIZE - 1) >> YUZU_PAGEBITS; for (u64 page = addr >> YUZU_PAGEBITS; page < page_end; ++page) { auto it = invalidation_cache.find(page); @@ -262,8 +274,22 @@ ShaderCache::Entry* ShaderCache::NewEntry(VAddr addr, VAddr addr_end, ShaderInfo } const ShaderInfo* ShaderCache::MakeShaderInfo(GenericEnvironment& env, VAddr cpu_addr) { + { + std::scoped_lock lock{rejected_mutex}; + if (const auto it = rejected_native.find(cpu_addr); it != rejected_native.end()) { + std::rethrow_exception(it->second); + } + } auto info = std::make_unique(); - if (const std::optional cached_hash{env.Analyze()}) { + std::optional analyzed; + try { + analyzed = env.Analyze(); + } catch (const Shader::Exception&) { + std::scoped_lock lock{rejected_mutex}; + rejected_native.emplace(cpu_addr, std::current_exception()); + throw; + } + if (const std::optional cached_hash{analyzed}) { info->unique_hash = *cached_hash; info->size_bytes = env.CachedSizeBytes(); } else { diff --git a/src/video_core/shader_cache.h b/src/video_core/shader_cache.h index 997f2b0371..81700ba0e6 100644 --- a/src/video_core/shader_cache.h +++ b/src/video_core/shader_cache.h @@ -7,6 +7,7 @@ #pragma once #include +#include #include #include #include @@ -152,6 +153,11 @@ private: ::Common::unordered_map> invalidation_cache; std::vector> storage; std::vector marked_for_removal; + + /// Native programs the lowering rejected, by address, so each draw or dispatch rethrows the + /// first error instead of lowering again; dropped when guest memory there is written. + std::mutex rejected_mutex; + ::Common::unordered_map rejected_native; }; } // namespace VideoCommon -- cgit v1.2.3 From 8f59920970674c47dcce8866d3e0659ef91a81b4 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 14:44:18 -0400 Subject: IMAD.SHL as a shift; EDEN_NX2_REJECT; log where rejected programs are bound - IMAD by a power of two lowers to SHL imm, the form Eden's jump-table tracking expects (a title 1 compute program failed "track indirect branch"). - EDEN_NX2_REJECT=: treat those native programs as rejected, to bisect a behaviour change without rebuilding. - A rejected graphics program logs its stage slot, header and slot registers once (for the task/mesh pipeline setup). - Notes: rows 87-92, title 1 status. --- NX2_NOTES.md | 15 ++++++++++++--- .../frontend/ampere/lower_to_maxwell.h | 10 ++++++++++ src/video_core/shader_cache.cpp | 18 +++++++++++++++++- src/video_core/shader_environment.cpp | 11 +++++++++++ 4 files changed, 50 insertions(+), 4 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/NX2_NOTES.md b/NX2_NOTES.md index d74a9cffb0..e623534dc8 100644 --- a/NX2_NOTES.md +++ b/NX2_NOTES.md @@ -142,6 +142,12 @@ With the patched decoder, games 1/3/4 only leave ~60-110 rare SIMD encodings in | 84 | title 1 GPU stalls (row 83): channels wait on semaphores 0x21f6093140/3040 until the 4 s timeout | NVN2 queue syncs run through a ring of QMDs (dependency 2 each: the previous entry's launch and macro 0xda, which patches release slot 52 and schedules/decrements via PCAS 3/7); macro 0xda reads the ring base from MME data RAM word 0x27, which each channel sets (0x21098d39, 0x213e8ff8, 0x213ea168) but Eden shared one data RAM for all channels, so every channel drove the same ring | per-channel MME data RAM, one QMD table for all channels (6c20e37); no stalls, title 1 reaches its title screen (world map with logo) | | 85 | title 1 address space | 214 = CreateSparseRegion, 215 = DestroySparseRegion (SDK symbols); SMMU space (16 GiB) ran out | 64 GiB device space (71b38ab), 215 handled (cc0c400) | | 86 | title 1 after its title screen: black frames | tens of thousands of rejected shaders: IPA interpolation modes, integer carry/third source, uniform comparison, formatted surface load/store, uniform constant load, unsupported opcodes | open | +| 87 | title 1 loading crash: Vulkan out of memory | an 8192x8192x68-layer virtual texture (~23 GB) | images over 2^28 texels are created sparse with no memory (`Unbacked`: no uploads/downloads) (96a9b81); the game reaches its real title screen ("Press L + R"), ground black (that texture has no data) | +| 88 | title 1 shader coverage | checked against nvdisasm (`nx2work/nvd.py` disassembles one instruction at a time, since a single illegal word stops nvdisasm); calls: RET may restore the link through IMAD.MOV with code between, library functions may EXIT, `@Pg EXIT Pc`; PLOP3/UPLOP3 any table, UP2UR, ULOP3 predicate output; UIMAD.WIDE, uniform USHF/FMNMX/I2F/F2I, LDS `[Ra.Xn+URb]`, ATOMG/RED/STG forms, PRMT, FSWZADD, REDUX.MAX | 1351c73, 52cda01, 49396b7: 1699 of 1712 dumped programs lower; left: 12 task/mesh programs (ISBERD/ISBEWR, row 89) and one SURED | +| 89 | task/mesh shaders | title 1 draws some geometry with NV-style task + mesh programs: the task program culls a batch and writes its payload with `ISBEWR.O.ATTR [Rn+0x10]` and the mesh count with `ISBEWR [0x4]`; mesh programs read it back with `ISBERD.ATTR.SKEW`. The most-drawn rejected program (0x114c459500) is a task program. Eden has no mesh pipeline | open: needs a task/mesh pipeline (VK_EXT_mesh_shader) or a compute + indirect-draw emulation | +| 90 | each rejected program was lowered and logged again on every draw (1.7M log lines in 13 min) | no negative cache | the shader cache keeps the first exception per address until that memory is written; logs deduplicated (bcb4ff0) | +| 91 | a compute program with a jump table: "Failed to track indirect branch" | Eden's tracker wants IMNMX imm -> SHL imm -> LDC -> BRX; IMAD.SHL was lowered to an XMAD multiply | IMAD by a power of two lowers to SHL (49396b7+) | +| 92 | with rows 88's compute forms, title 1 stays on its loading map (file reads stop ~46 s in) | one of the eight newly lowered compute programs (an FFT with LDS, ATOMG, RED, STG.U16, REDUX.MAX) | open: bisect with `EDEN_NX2_REJECT=` (rejects those programs without a rebuild) | Run settings used (in `build/bin/user/config/sdl2-config.ini`): `memory_layout_mode=4` (12 GB). Debug knobs: `EDEN_NX2_AUTO_STUB`, `EDEN_NX2_IPC_TRACE`, `EDEN_NX2_SVC_TRACE`, @@ -166,9 +172,12 @@ Scripted input: `EDEN_NX2_KEY_SCRIPT="50:4,50.5:-4,55:79,55.4:-79,57:4,57.5:-4"` **Title 3**: last checked before the graphics work; should now track title 4. -**Title 1** (`first switch game\exefs`, 0400C3F00006E000, `Drive.nss`): reached its main loop and -the empty-command-buffer stop that blocker 45 fixed for both titles; not retested since the -macro/shader/present work. +**Title 1** (`first switch game\exefs`, 0400C3F00006E000, `Drive.nss`): boots, shows the world-map +loading screen and reaches its title screen ("Press L + R to start") with sky, buildings, props, +kart and logo drawn. Wrong: magenta/cyan colour cast, black ground (unbacked virtual texture, row +87), task/mesh geometry missing (row 89), ~10-17 fps. Launch with frame capture: +`EDEN_NX2_FRAME_CAPTURE=C:/dev/frames_t1x EDEN_NX2_AUTO_STUB=1 EDEN_NX2_SHADER_DUMP=C:/dev/shaders_title1 ./eden-cli.exe -g "C:/switch/first switch game/exefs/main"`. +Offline shader check: `nx2work/lowertest/lowertest.exe C:/dev/shaders_title1 [-all [op]]`. **Title 2**: main/sdk/subsdk0 fail LZ4 and SHA-256 checks; can't be run from these files. diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index c957bc29b6..0bcb6a737a 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -547,6 +547,16 @@ inline LoweredProgram LowerToMaxwell(std::span code) { (u64{insn.Bits(64, 8)} << 20) | (u64{negate_c} << 48)); break; } + const u32 multiplier = insn.words[1]; + if (insn.Bits(64, 8) == 255 && !negate_c && multiplier > 1 && !(multiplier & (multiplier - 1)) && + (insn.words[2] & ~0x2ffU) == 0x078e0000) { + // IMAD.SHL by a power of two: one shift (also the form Eden's jump-table tracking + // expects between the clamp and the table load). + u32 shift = 0; + while ((1U << shift) != multiplier) ++shift; + emit(base(0x3848000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{shift} << 20)); + break; + } if (insn.words[1] != 48 || insn.words[2] != 0x078e02ff) { if ((insn.words[2] & ~0xaffU) != 0x078e0000) return fail(index, "native integer multiply mode"); diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index f92fe02068..15c1adb7c9 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -5,6 +5,7 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include #include #include #include @@ -81,7 +82,22 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) try { if (!shader_info) { const u32 start_address{native ? 0 : shader_config.offset}; GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, native ? shader_addr : base_addr, start_address}; - shader_info = MakeShaderInfo(env, *cpu_shader_addr); + try { + shader_info = MakeShaderInfo(env, *cpu_shader_addr); + } catch (const Shader::Exception&) { + // Once per program: which stage slot it is bound to, its header and the slot's + // registers (to tell how a pipeline with task/mesh programs is set up). + static std::unordered_set described; + if (native && described.insert(shader_addr).second) { + std::array header{}; + gpu_memory->ReadBlockUnsafe(shader_addr, header.data(), sizeof(header)); + std::array config{}; + std::memcpy(config.data(), &shader_config, sizeof(config)); + LOG_ERROR(HW_GPU, "Rejected program {:#x} is bound to stage {}: header {:08x}, slot {:08x}", + shader_addr, index, fmt::join(header, " "), fmt::join(config, " ")); + } + throw; + } } shader_infos[index] = shader_info; unique_hashes[index] = shader_info->unique_hash; diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp index 5a14211073..96306d9a27 100644 --- a/src/video_core/shader_environment.cpp +++ b/src/video_core/shader_environment.cpp @@ -10,6 +10,7 @@ #include #include #include +#include #include #include #include @@ -354,6 +355,16 @@ void GenericEnvironment::InitializeNativeCode(u32 source_header_size) { if (status == Shader::Ampere::ProgramCaptureState::Result::Complete) break; } auto lowered = Shader::Ampere::LowerToMaxwell(native); + // EDEN_NX2_REJECT=,...: treat these programs as rejected (to find which newly + // lowered program changes a game's behaviour without rebuilding). + if (static const std::string reject_list = [] { + const char* text = std::getenv("EDEN_NX2_REJECT"); + return text ? "," + std::string{text} + "," : std::string{}; + }(); + !reject_list.empty() && + reject_list.find(fmt::format(",{:x},", program_base)) != std::string::npos) { + throw Shader::RuntimeError("Native shader at {:#x} rejected by EDEN_NX2_REJECT", program_base); + } // EDEN_NX2_SHADER_DUMP=: keep every captured native program (named by its key) so the // lowering can be tested offline against all of a game's shaders. if (static const char* const dump_directory = std::getenv("EDEN_NX2_SHADER_DUMP"); -- cgit v1.2.3 From 0a692577b5c7d63a252d5099e39c088a117e7818 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 14:50:32 -0400 Subject: Jump-table programs end at their self-branch; aligned wide zero stores - With a BRX the jump targets live in a constant buffer, so capture and lowering keep going past EXITs to the compiler's closing branch-to-self (a title 1 compute program jumped to 0xae0, just past the first EXIT). - 128-bit stores of RZ use R248-R251 (Eden wants 4-aligned data registers). - The rejected-program cache covers environment construction, where native programs are lowered (it never applied before), and logs the stage slot. Title 1 reaches its title screen with the ground drawn. --- .../frontend/ampere/lower_to_maxwell.h | 17 ++++++--- .../frontend/ampere/program_capture.h | 8 ++++ src/video_core/shader_cache.cpp | 43 +++++++++++++--------- src/video_core/shader_cache.h | 5 +++ 4 files changed, 49 insertions(+), 24 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 0bcb6a737a..0cebaa4363 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -50,6 +50,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { std::vector call_targets; std::vector call_links; bool main_exited = false; + bool computed_branch = false; // a jump table: the program ends at its self-branch, not an EXIT bool uses_uniform_bank = false; for (const auto& instruction : code) { const u32 op = instruction.Bits(0, 12); @@ -1264,9 +1265,10 @@ inline LoweredProgram LowerToMaxwell(std::span code) { data + (size < 4 ? 1U : 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; + // Zeros wider than one register: store from zeroed scratch registers (aligned to the + // access: R248-R251 for 128 bits, R250-R251 for 64). + data = size == 6 ? 248 : 250; + for (u32 i = 0; i < (1U << (size - 4)); ++i) mov_imm(data + i, 0); } u32 address = insn.Bits(24, 8); bool extended = insn.Bits(72, 1) != 0; @@ -2090,6 +2092,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { if (source == ~0U) return fail(index, "native uniform computed branch source"); emit(0xe25000000000000fULL | predicate | (u64{source} << 8) | ((static_cast(offset) & 0xffffff) << 20)); + computed_branch = true; break; } case 0x944: { // Verified CALL.REL.NOINC with R55 holding the return PC @@ -2149,7 +2152,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { if (s64{index} * 16 + 16 + relative != 0) return fail(index, "native return relative base"); emit(0xe32000000007000fULL); - bool complete = true; + bool complete = !computed_branch; for (u32 target : call_targets) complete &= target <= index; if (complete) return finish(index); break; @@ -2177,6 +2180,8 @@ inline LoweredProgram LowerToMaxwell(std::span code) { if (target < 0 || target % 16) return fail(index, "native branch destination alignment"); relocations.push_back({static_cast(result.instructions.size()), target / 16}); emit(0xe24000000000000fULL | branch_predicate); + if (computed_branch && target == s64{index} * 16 && branch_predicate == (u64{7} << 16)) + return finish(index); // the closing self-branch break; } case 0x94d: { // EXIT @@ -2198,7 +2203,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { if (call_targets.empty()) { const bool pending_branch = std::any_of(relocations.begin(), relocations.end(), [&](const Relocation& relocation) { return relocation.target > index; }); - if (!pending_branch) return finish(index); + if (!pending_branch && !computed_branch) return finish(index); break; } if (main_exited) { @@ -2207,7 +2212,7 @@ inline LoweredProgram LowerToMaxwell(std::span code) { bool complete = std::none_of(relocations.begin(), relocations.end(), [&](const Relocation& relocation) { return relocation.target > index; }); for (u32 target : call_targets) complete &= target <= index; - if (complete) return finish(index); + if (complete && !computed_branch) return finish(index); break; } for (u32 target : call_targets) { diff --git a/src/shader_recompiler/frontend/ampere/program_capture.h b/src/shader_recompiler/frontend/ampere/program_capture.h index f09bf9735d..7e308de4ab 100644 --- a/src/shader_recompiler/frontend/ampere/program_capture.h +++ b/src/shader_recompiler/frontend/ampere/program_capture.h @@ -27,7 +27,14 @@ public: } else { last_forward_target = std::max(last_forward_target, static_cast(target / 16)); } + // The compiler ends every program with an unconditional branch to itself; with a jump + // table (whose targets live in a constant buffer) that is the only reliable end. + if (opcode == 0x947 && computed_branch && instruction.Bits(12, 4) == 7 && + target == s64{index} * 16) + return Result::Complete; } + if (opcode == 0x949 || opcode == 0x958) computed_branch = true; + if (computed_branch) return Result::Continue; if (opcode == 0x94d && instruction.Bits(12, 4) == 7 && index >= last_forward_target) { if (!last_call_target) return Result::Complete; main_exited = true; @@ -42,5 +49,6 @@ private: u32 last_call_target{}; u32 last_forward_target{}; bool main_exited{}; + bool computed_branch{}; }; } // namespace Shader::Ampere diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index 15c1adb7c9..da5417be02 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -81,10 +81,13 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) try { const ShaderInfo* shader_info{TryGet(*cpu_shader_addr)}; if (!shader_info) { const u32 start_address{native ? 0 : shader_config.offset}; - GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, native ? shader_addr : base_addr, start_address}; try { + // Native programs are lowered while the environment is built. + RethrowIfRejected(*cpu_shader_addr); + GraphicsEnvironment env{*maxwell3d, *gpu_memory, program, native ? shader_addr : base_addr, start_address}; shader_info = MakeShaderInfo(env, *cpu_shader_addr); } catch (const Shader::Exception&) { + RememberRejection(*cpu_shader_addr); // Once per program: which stage slot it is bound to, its header and the slot's // registers (to tell how a pipeline with task/mesh programs is set up). static std::unordered_set described; @@ -145,8 +148,26 @@ const ShaderInfo* ShaderCache::ComputeShader() { if (const ShaderInfo* const shader = TryGet(*cpu_shader_addr)) { return shader; } - ComputeEnvironment env{*kepler_compute, *gpu_memory, program_base, qmd.program_start}; - return MakeShaderInfo(env, *cpu_shader_addr); + try { + RethrowIfRejected(*cpu_shader_addr); + ComputeEnvironment env{*kepler_compute, *gpu_memory, program_base, qmd.program_start}; + return MakeShaderInfo(env, *cpu_shader_addr); + } catch (const Shader::Exception&) { + RememberRejection(*cpu_shader_addr); + throw; + } +} + +void ShaderCache::RethrowIfRejected(VAddr addr) { + std::scoped_lock lock{rejected_mutex}; + if (const auto it = rejected_native.find(addr); it != rejected_native.end()) { + std::rethrow_exception(it->second); + } +} + +void ShaderCache::RememberRejection(VAddr addr) { + std::scoped_lock lock{rejected_mutex}; + rejected_native.emplace(addr, std::current_exception()); } void ShaderCache::GetGraphicsEnvironments(GraphicsEnvironments& result, @@ -290,22 +311,8 @@ ShaderCache::Entry* ShaderCache::NewEntry(VAddr addr, VAddr addr_end, ShaderInfo } const ShaderInfo* ShaderCache::MakeShaderInfo(GenericEnvironment& env, VAddr cpu_addr) { - { - std::scoped_lock lock{rejected_mutex}; - if (const auto it = rejected_native.find(cpu_addr); it != rejected_native.end()) { - std::rethrow_exception(it->second); - } - } auto info = std::make_unique(); - std::optional analyzed; - try { - analyzed = env.Analyze(); - } catch (const Shader::Exception&) { - std::scoped_lock lock{rejected_mutex}; - rejected_native.emplace(cpu_addr, std::current_exception()); - throw; - } - if (const std::optional cached_hash{analyzed}) { + if (const std::optional cached_hash{env.Analyze()}) { info->unique_hash = *cached_hash; info->size_bytes = env.CachedSizeBytes(); } else { diff --git a/src/video_core/shader_cache.h b/src/video_core/shader_cache.h index 81700ba0e6..eac2c5ba75 100644 --- a/src/video_core/shader_cache.h +++ b/src/video_core/shader_cache.h @@ -144,6 +144,11 @@ private: /// @brief Create a new shader entry and register it const ShaderInfo* MakeShaderInfo(GenericEnvironment& env, VAddr cpu_addr); + /// Rethrows the error a native program at this address was rejected with, if any. + void RethrowIfRejected(VAddr addr); + /// Remembers the exception being handled as this address's rejection. + void RememberRejection(VAddr addr); + Tegra::MaxwellDeviceMemoryManager& device_memory; mutable std::mutex lookup_mutex; -- cgit v1.2.3 From 25b674eabdb92212a6d8149d6616b2af59317695 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Wed, 28 Oct 2026 23:07:49 -0400 Subject: nx2: measure the TSC on AMD, no rethrows for rejected shaders, debug hooks off the hot path - AMD CPUs have no CPUID leaf 0x15, so the TSC frequency stayed 0 and every guest CNTPCT read went through QueryPerformanceCounter (~10% of the game's main core in title 1). Measure it like Intel CPUs without a crystal frequency. - A rejected native shader was rethrown (two C++ throws) every time it was bound; return false instead (~1-2% of the GPU thread). - The semaphore watch, pushbuffer entry hashing, first-write trace and per-macro debug logs ran on every command word or macro call; they now cost a flag check unless their environment variable is set (~4% of the GPU thread). Also starts the macro engine's profile zone. --- src/common/cpu_features.cpp | 7 +++++++ src/video_core/dma_pusher.cpp | 28 ++++++++++++++++++++-------- src/video_core/engines/maxwell_3d.cpp | 9 ++++++--- src/video_core/macro.cpp | 9 +++++++-- src/video_core/shader_cache.cpp | 12 ++++++++++++ src/video_core/shader_cache.h | 2 ++ 6 files changed, 54 insertions(+), 13 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/common/cpu_features.cpp b/src/common/cpu_features.cpp index 8ccb6847f7..52fd003c44 100644 --- a/src/common/cpu_features.cpp +++ b/src/common/cpu_features.cpp @@ -225,6 +225,13 @@ const CPUCaps g_cpu_caps = [] { } } + // AMD CPUs have no leaf 0x15 (Zen 4 stops at 0x10): measure their invariant TSC too, or every + // guest clock read (CNTPCT, which games read thousands of times a frame) goes through the + // host's much slower QueryPerformanceCounter. + if (caps.invariant_tsc && caps.tsc_frequency == 0) { + caps.tsc_frequency = X64::EstimateRDTSCFrequency(); + } + if (max_std_fn >= 0x16) { __cpuid(cpu_id, 0x16); caps.base_frequency = cpu_id[0]; diff --git a/src/video_core/dma_pusher.cpp b/src/video_core/dma_pusher.cpp index a1d9ef3f91..3cc5afe967 100644 --- a/src/video_core/dma_pusher.cpp +++ b/src/video_core/dma_pusher.cpp @@ -299,19 +299,26 @@ void DmaPusher::SetState(const CommandHeader& command_header) { // 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). +static const u32 g_semaphore_watch = [] { + const char* value = std::getenv("EDEN_NX2_SEM_WATCH"); + return value ? static_cast(std::strtoull(value, nullptr, 16)) : 0U; +}(); + +// EDEN_NX2_PB_CHECK also keeps the last parsed entries, for DumpCurrentEntry. +static const bool g_record_parsed_entries = std::getenv("EDEN_NX2_PB_CHECK") != nullptr; + 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) { + const u32 watch = g_semaphore_watch; + if ((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 (g_semaphore_watch != 0) [[unlikely]] { + WatchArgument(dma_state.method, argument); + } if (dma_state.method < non_puller_methods) { puller.CallPullerMethod(*this, Engines::Puller::MethodCall{ dma_state.method, @@ -336,8 +343,10 @@ 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 (g_semaphore_watch != 0) [[unlikely]] { + 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); @@ -382,6 +391,9 @@ void DmaPusher::CheckSubmitHash(const CommandListHeader& header, u64 submit_hash void DmaPusher::RecordParsedEntry(const CommandListHeader& header, std::span commands) { + if (!g_record_parsed_entries) { + return; + } u64 hash = 0xcbf29ce484222325ULL; for (const CommandHeader& command : commands) { hash = (hash ^ command.argument) * 0x100000001b3ULL; diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp index 7884d8009c..60efef9b8e 100644 --- a/src/video_core/engines/maxwell_3d.cpp +++ b/src/video_core/engines/maxwell_3d.cpp @@ -393,10 +393,11 @@ void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, const std::vec } // EDEN_NX2_3D_METHODS: log the first write to each 3D register (finds methods Eden doesn't model). +static const bool g_trace_first_method_writes = std::getenv("EDEN_NX2_3D_METHODS") != nullptr; + 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]) { + if (method < seen.size() && !seen[method]) { seen[method] = true; LOG_INFO(HW_GPU, "3D method first write: {:#x} (byte {:#x}) = {:#x}", method, method * 4, argument); } @@ -419,7 +420,9 @@ 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); + if (g_trace_first_method_writes) [[unlikely]] { + 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); diff --git a/src/video_core/macro.cpp b/src/video_core/macro.cpp index d3ba353ca0..c74cd6087a 100644 --- a/src/video_core/macro.cpp +++ b/src/video_core/macro.cpp @@ -36,6 +36,7 @@ #include "common/bit_field.h" #include "common/logging.h" #include "video_core/engines/maxwell_3d.h" +#include "video_core/nx2_profile.h" namespace Tegra { @@ -1343,6 +1344,7 @@ static void Dump(u64 hash, std::span code, bool decompiled = false) { void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u32 method, std::span parameters) { + VideoCommon::Nx2Profile::Scope profile_scope{VideoCommon::Nx2Profile::Zone::Macro}; if (native_turing) { maxwell3d.RefreshParameters(); const auto context = reinterpret_cast(this); @@ -1352,8 +1354,11 @@ void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u std::ofstream dump("native-graphics-mme-" + std::to_string(context) + ".bin", std::ios::binary); dump.write(reinterpret_cast(native_code.data()), sizeof(native_code)); } - 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, ",")); + // EDEN_NX2_MME_LOG: log every macro call (runs ~750,000 times a second in title 1). + if (static const bool log_calls = std::getenv("EDEN_NX2_MME_LOG") != nullptr; log_calls) { + 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, ",")); + } // Reused between calls (a macro never calls another: methods >= 0xe00 are refused below). auto& fifo = native_fifo; fifo.assign(parameters.begin(), parameters.end()); diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index da5417be02..73a2bf6646 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -79,6 +79,13 @@ bool ShaderCache::RefreshStages(std::array& unique_hashes) try { return false; } const ShaderInfo* shader_info{TryGet(*cpu_shader_addr)}; + if (!shader_info && native && IsRejected(*cpu_shader_addr)) { + // A rejected program fails the same way every time it is bound: skip its draws + // without rethrowing (two C++ throws each time cost ~1-2% of the GPU thread). + maxwell3d->GetNativeMacroEngine().native_execution_failed = true; + last_shaders_valid = false; + return false; + } if (!shader_info) { const u32 start_address{native ? 0 : shader_config.offset}; try { @@ -165,6 +172,11 @@ void ShaderCache::RethrowIfRejected(VAddr addr) { } } +bool ShaderCache::IsRejected(VAddr addr) { + std::scoped_lock lock{rejected_mutex}; + return rejected_native.contains(addr); +} + void ShaderCache::RememberRejection(VAddr addr) { std::scoped_lock lock{rejected_mutex}; rejected_native.emplace(addr, std::current_exception()); diff --git a/src/video_core/shader_cache.h b/src/video_core/shader_cache.h index eac2c5ba75..9ffb7bb44d 100644 --- a/src/video_core/shader_cache.h +++ b/src/video_core/shader_cache.h @@ -146,6 +146,8 @@ private: /// Rethrows the error a native program at this address was rejected with, if any. void RethrowIfRejected(VAddr addr); + /// Whether a native program at this address was rejected. + bool IsRejected(VAddr addr); /// Remembers the exception being handled as this address's rejection. void RememberRejection(VAddr addr); -- cgit v1.2.3 From d56a30137df183de548ac49151ddeeecc1ab0f16 Mon Sep 17 00:00:00 2001 From: Takoroka <137927095+Takoroka@users.noreply.github.com> Date: Thu, 29 Oct 2026 00:45:20 -0400 Subject: nx2: memo for compute macros; one lock per inline upload; cheaper shader invalidation - Compute engine macros replay repeated calls like 3D ones (title 1's compute macro path fell from ~27% to ~14% of the GPU thread). - AccelerateInlineToMemory takes the buffer and texture cache mutexes together instead of one after the other. - ShaderCache::InvalidatePagesInRegion looks at the rejected programs only when the write can reach one (bounds kept as they are added); it ran a locked scan for every inline upload and pre-draw invalidation. Inline uploads went from ~10% to ~4% of the GPU thread. --- src/video_core/engines/kepler_compute.cpp | 55 ++++++++++++++++++++---- src/video_core/engines/kepler_compute.h | 1 + src/video_core/renderer_vulkan/vk_rasterizer.cpp | 6 +-- src/video_core/shader_cache.cpp | 9 +++- src/video_core/shader_cache.h | 4 ++ 5 files changed, 62 insertions(+), 13 deletions(-) (limited to 'src/video_core/shader_cache.cpp') diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 39c033d989..7792bc9b5d 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -147,18 +147,57 @@ bool KeplerCompute::ProcessNativeMethod(Core::System& system, u32 method, u32 ar } if (native_data_ram) native_macro->native_interpreter.active_dram = native_data_ram; const u32 dram0_before = (*native_macro->native_interpreter.active_dram)[0x27]; + const u32 start = native_macro_positions[entry]; + const auto emit = [&](u32 address, u32 value) { + static u32 diagnostic_writes{}; + 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); + return true; + }; + // Repeated calls replay their writes, as for 3D macros (macro_memo.h). + static const bool use_memo = [] { + const char* value = std::getenv("EDEN_NX2_MME_MEMO"); + return value == nullptr || value[0] != '0'; + }(); + auto& memo = native_macro->native_memo; + if (use_memo && !watched) { + memo.Synchronize(native_macro->native_code_generation); + if (std::remove_reference_t::Effects effects; + memo.Find(start, native_parameters, *native_macro->native_interpreter.active_dram, + effects)) { + auto& data_ram = *native_macro->native_interpreter.active_dram; + for (size_t i = 0; i < effects.data_ram_writes.size(); i += 2) { + data_ram[effects.data_ram_writes[i]] = effects.data_ram_writes[i + 1]; + } + for (size_t i = 0; i < effects.writes.size(); i += 2) { + emit(effects.writes[i], effects.writes[i + 1]); + } + native_parameters.clear(); + return true; + } + native_macro->native_writes.clear(); + } + // The interpreter consumes native_parameters as its FIFO: keep the call's own parameters. + native_memo_parameters.assign(native_parameters.begin(), native_parameters.end()); + const u64 code_generation = native_macro->native_code_generation; const auto result = native_macro->native_interpreter.Execute(native_macro->native_code, - native_macro_positions[entry], native_parameters, + start, native_parameters, [&](u32 address) { return address < native_registers.size() ? native_registers[address] : 0U; }, [&](u32 address, u32 value) { - static u32 diagnostic_writes{}; - 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); - return true; + if (use_memo) { + native_macro->native_writes.emplace_back(address, value); + } + return emit(address, value); }); + const auto& interpreter = native_macro->native_interpreter; + if (use_memo && !watched && result.success && !interpreter.read_registers && + !interpreter.reread_data_ram && code_generation == native_macro->native_code_generation) { + memo.Store(start, native_memo_parameters, interpreter.data_ram_reads, + interpreter.data_ram_writes, native_macro->native_writes); + } if (static const bool mme_dump = std::getenv("EDEN_NX2_MME_DUMP") != nullptr; mme_dump && (*native_macro->native_interpreter.active_dram)[0x27] != dram0_before) { LOG_INFO(HW_GPU, "MME data RAM[0x27] {:#x} -> {:#x} by engine {} macro start={:#x}", diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index 7cecab378b..5ebf8edf38 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -254,6 +254,7 @@ private: std::array native_registers{}; std::span native_macro_positions; std::vector native_parameters; + std::vector native_memo_parameters; ///< The current call's parameters, for the memo std::optional native_executing_macro; u32 native_bind_pointer{}; diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index decfb0d136..6bd267e566 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -1251,13 +1251,11 @@ void RasterizerVulkan::AccelerateInlineToMemory(GPUVAddr address, size_t copy_si } gpu_memory->WriteBlockUnsafe(address, memory.data(), copy_size); { - std::unique_lock lock{buffer_cache.mutex}; + // Both caches under one lock (title 1 makes thousands of inline uploads a second). + std::scoped_lock lock{buffer_cache.mutex, texture_cache.mutex}; if (!buffer_cache.InlineMemory(*cpu_addr, copy_size, memory)) { buffer_cache.WriteMemory(*cpu_addr, copy_size); } - } - { - std::scoped_lock lock_texture{texture_cache.mutex}; texture_cache.WriteMemory(*cpu_addr, copy_size); } pipeline_cache.InvalidateRegion(*cpu_addr, copy_size); diff --git a/src/video_core/shader_cache.cpp b/src/video_core/shader_cache.cpp index 73a2bf6646..dfe80abac3 100644 --- a/src/video_core/shader_cache.cpp +++ b/src/video_core/shader_cache.cpp @@ -180,6 +180,10 @@ bool ShaderCache::IsRejected(VAddr addr) { void ShaderCache::RememberRejection(VAddr addr) { std::scoped_lock lock{rejected_mutex}; rejected_native.emplace(addr, std::current_exception()); + rejected_lowest.store((std::min)(rejected_lowest.load(std::memory_order_relaxed), addr), + std::memory_order_relaxed); + rejected_highest.store((std::max)(rejected_highest.load(std::memory_order_relaxed), addr), + std::memory_order_relaxed); } void ShaderCache::GetGraphicsEnvironments(GraphicsEnvironments& result, @@ -232,7 +236,10 @@ void ShaderCache::Register(std::unique_ptr data, VAddr addr, size_t void ShaderCache::InvalidatePagesInRegion(VAddr addr, size_t size) { const VAddr addr_end = addr + size; - { + // Rejected programs are few and in a few places: look at them only for writes that can + // reach them (this runs for every inline upload and pre-draw invalidation). + if (addr_end > rejected_lowest.load(std::memory_order_relaxed) && + addr <= rejected_highest.load(std::memory_order_relaxed)) { std::scoped_lock lock{rejected_mutex}; boost::unordered::erase_if(rejected_native, [&](const auto& rejected) { return rejected.first >= addr && rejected.first < addr_end; diff --git a/src/video_core/shader_cache.h b/src/video_core/shader_cache.h index 9ffb7bb44d..f4179dd435 100644 --- a/src/video_core/shader_cache.h +++ b/src/video_core/shader_cache.h @@ -7,6 +7,7 @@ #pragma once #include +#include #include #include #include @@ -165,6 +166,9 @@ private: /// first error instead of lowering again; dropped when guest memory there is written. std::mutex rejected_mutex; ::Common::unordered_map rejected_native; + /// Bounds of every address ever rejected (never shrinks): a quick test for invalidations. + std::atomic rejected_lowest{~VAddr{}}; + std::atomic rejected_highest{}; }; } // namespace VideoCommon -- cgit v1.2.3