Merge Switch 2 work from local Eden (nx2 at b0dd3fe)

1 files changed, 203 insertions, 3 deletions

diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp
index 4d5db7e511..96306d9a27 100644
--- a/src/video_core/shader_environment.cpp
+++ b/src/video_core/shader_environment.cpp
@@ -5,7 +5,12 @@
 // SPDX-License-Identifier: GPL-2.0-or-later
 
 #include <algorithm>
+#include <chrono>
+#include <cstdlib>
+#include <thread>
 #include <filesystem>
+#include <cstring>
+#include <string>
 #include <fstream>
 #include <memory>
 #include <optional>
@@ -21,6 +26,11 @@
 #include "common/settings.h"
 #include <ranges>
 #include "shader_recompiler/environment.h"
+#include "shader_recompiler/exception.h"
+#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h"
+#include "shader_recompiler/frontend/ampere/program_override.h"
+#include "shader_recompiler/frontend/ampere/program_capture.h"
+#include "shader_recompiler/frontend/ampere/program_address_map.h"
 #include "video_core/engines/kepler_compute.h"
 #include "video_core/memory_manager.h"
 #include "video_core/shader_environment.h"
@@ -149,11 +159,26 @@ u64 GenericEnvironment::ReadInstruction(u32 address) {
     if (address >= cached_lowest && address < cached_highest) {
         return code[(address - cached_lowest) / INST_SIZE];
     }
+    if (native_source_size) {
+        throw Shader::RuntimeError("Lowered native instruction address out of range: {:#x}", address);
+    }
     has_unbound_instructions = true;
     return gpu_memory->Read<u64>(program_base + address);
 }
 
+u32 GenericEnvironment::TranslateProgramAddress(u32 address) const {
+    if (const auto translated = Shader::Ampere::NativeProgramAddress(code, address)) return *translated;
+    if (native_source_size || Shader::Ampere::HasNativeAddressMap(code))
+        throw Shader::RuntimeError("Invalid native indirect branch address {:#x}", address);
+    return address;
+}
+
 std::optional<u64> GenericEnvironment::Analyze() {
+    if (native_source_size) {
+        // Include the converted header and a version tag to isolate this experimental ABI.
+        return Common::CityHash64(reinterpret_cast<const char*>(code.data()),
+                                  code.size() * sizeof(u64)) ^ 0x4f554e4345000002ULL;
+    }
     const std::optional<u64> size{TryFindSize()};
     if (!size) {
         return std::nullopt;
@@ -164,6 +189,12 @@ std::optional<u64> GenericEnvironment::Analyze() {
 }
 
 void GenericEnvironment::SetCachedSize(size_t size_bytes) {
+    if (native_source_size) {
+        if (size_bytes != CachedSizeBytes()) {
+            throw Shader::RuntimeError("Native shader changed while constructing pipeline");
+        }
+        return;
+    }
     cached_lowest = start_address;
     cached_highest = start_address + static_cast<u32>(size_bytes);
     code.resize(CachedSizeWords());
@@ -291,6 +322,146 @@ Tegra::Texture::TICEntry GenericEnvironment::ReadTextureInfo(GPUVAddr tic_addr,
     return entry;
 }
 
+void GenericEnvironment::InitializeNativeCode(u32 source_header_size) {
+    using Shader::Ampere::Instruction;
+    const u32 native_header_size = source_header_size;
+    constexpr u32 maximum_instructions = Shader::Ampere::ProgramCaptureState::MaximumInstructions;
+    std::vector<Instruction> native;
+    Shader::Ampere::ProgramCaptureState capture;
+    for (u32 index = 0; index < maximum_instructions; ++index) {
+        const GPUVAddr address = program_base + native_header_size + index * sizeof(Instruction);
+        if (!gpu_memory->IsFullyMappedRange(address, sizeof(Instruction))) {
+            // The program may end before a branch target past its last EXIT (the capture can't
+            // tell); stop at the mapping's end and let the lowering judge what was captured.
+            if (!native.empty()) break;
+            throw Shader::RuntimeError("Unmapped native shader instruction at {:#x}", address);
+        }
+        auto& instruction = native.emplace_back();
+        gpu_memory->ReadBlock(address, &instruction, sizeof(instruction));
+        const auto status = capture.Observe(instruction, index);
+        if (status == Shader::Ampere::ProgramCaptureState::Result::InvalidTarget) {
+            if (const char* dir = std::getenv("EDEN_NX2_SHADER_DUMP")) {
+                const auto path = fmt::format("{}/rejected-{:x}.bin", dir, program_base + native_header_size);
+                if (std::FILE* file = std::fopen(path.c_str(), "wb")) {
+                    std::fwrite(native.data(), sizeof(Instruction), std::min<size_t>(native.size(), 4096), file);
+                    std::fclose(file);
+                }
+            }
+            throw Shader::RuntimeError("Native shader branch outside capture bounds at {:#x} +{:#x} [{:08x} {:08x} {:08x} {:08x}]",
+                                       program_base + native_header_size, index * 16,
+                                       instruction.words[0], instruction.words[1],
+                                       instruction.words[2], instruction.words[3]);
+        }
+        if (status == Shader::Ampere::ProgramCaptureState::Result::Complete) break;
+    }
+    auto lowered = Shader::Ampere::LowerToMaxwell(native);
+    // EDEN_NX2_REJECT=<hex address>,...: 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=<dir>: 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");
+        dump_directory != nullptr) {
+        const auto folder = std::filesystem::path{dump_directory};
+        const auto path = folder / fmt::format("{:016x}-{}-{:x}.bin", Shader::Ampere::NativeProgramKey(native),
+                                               source_header_size == 0 ? "cs" : "gfx", program_base);
+        if (!std::filesystem::exists(path)) {
+            std::filesystem::create_directories(folder);
+            std::ofstream dump(path, std::ios::binary);
+            dump.write(reinterpret_cast<const char*>(native.data()),
+                       native.size() * sizeof(Shader::Ampere::Instruction));
+        }
+    }
+    if (!lowered) {
+        if (const char* directory = std::getenv("EDEN_NATIVE_SHADER_OVERRIDE_DIR")) {
+            const auto key = fmt::format("{:016x}", Shader::Ampere::NativeProgramKey(native));
+            const auto folder = std::filesystem::path{directory};
+            std::filesystem::create_directories(folder);
+            const auto source_path = folder / (key + ".source.bin");
+            const auto override_path = folder / (key + ".override.bin");
+            {
+                std::ofstream dump(source_path, std::ios::binary);
+                dump.write(reinterpret_cast<const char*>(native.data()), native.size() * sizeof(Shader::Ampere::Instruction));
+            }
+            {
+                // Retain mapped bytes beyond the first EXIT for investigating
+                // native CALL targets in compiler library subroutines.
+                std::ofstream mapped(folder / (key + ".mapped.bin"), std::ios::binary);
+                std::array<u8, 0x1000> block{};
+                for (u32 offset = 0; offset < 0x10000; offset += 0x1000) {
+                    if (!gpu_memory->IsFullyMappedRange(program_base + offset, block.size())) break;
+                    gpu_memory->ReadBlock(program_base + offset, block.data(), block.size());
+                    mapped.write(reinterpret_cast<const char*>(block.data()), block.size());
+                }
+            }
+            LOG_WARNING(Shader, "Native shader paused: key={}, offset={:#x}, reason={}, source={}",
+                        key, lowered.error_offset, lowered.error, source_path.string());
+            const auto deadline = std::chrono::steady_clock::now() + std::chrono::minutes{30};
+            while (!gpu_memory->IsShuttingDown() && std::chrono::steady_clock::now() < deadline) {
+                if (std::filesystem::exists(override_path)) {
+                    std::ifstream file(override_path, std::ios::binary | std::ios::ate);
+                    const auto length = file.tellg();
+                    if (length <= 0 || length > 16 * 1024 * 1024) {
+                        throw Shader::RuntimeError("Invalid native shader override size");
+                    }
+                    std::vector<u8> bytes(static_cast<size_t>(length));
+                    file.seekg(0);
+                    file.read(reinterpret_cast<char*>(bytes.data()), bytes.size());
+                    if (!file) throw Shader::RuntimeError("Native shader override read failed");
+                    const u32 source_bytes = Shader::Ampere::ProgramOverrideSourceSize(bytes);
+                    if (source_bytes > native.size() * sizeof(Shader::Ampere::Instruction)) {
+                        if (!gpu_memory->IsFullyMappedRange(program_base + native_header_size, source_bytes))
+                            throw Shader::RuntimeError("Unmapped extended native shader source");
+                        native.resize(source_bytes / sizeof(Shader::Ampere::Instruction));
+                        gpu_memory->ReadBlock(program_base + native_header_size, native.data(), source_bytes);
+                    }
+                    lowered = Shader::Ampere::DecodeProgramOverride(bytes, native);
+                    if (!lowered) throw Shader::RuntimeError("Native shader override rejected: {}", lowered.error);
+                    LOG_WARNING(Shader, "Native shader resumed: key={}, instructions={}", key, lowered.instructions.size());
+                    break;
+                }
+                std::this_thread::sleep_for(std::chrono::milliseconds{100});
+            }
+        }
+    }
+    if (!lowered) {
+        const size_t failed = lowered.error_offset / sizeof(Shader::Ampere::Instruction);
+        const auto words = failed < native.size() ? native[failed].words : std::array<u32, 4>{};
+        throw Shader::RuntimeError(
+            "Native shader at {:#x}, instruction {:#x}: {} [op {:03x} {:08x} {:08x} {:08x} {:08x}]",
+            program_base, lowered.error_offset, lowered.error, words[0] & 0xfff, words[0],
+            words[1], words[2], words[3]);
+    }
+    if (lowered.required_texture_buffer != ~0U) {
+        // Native TEX carries its cbuf index explicitly; TEXS obtains it from the environment.
+        texture_bound = lowered.required_texture_buffer;
+    }
+    native_source_size = native_header_size + lowered.source_size;
+    code.resize(sizeof(sph) / sizeof(u64));
+    std::memcpy(code.data(), &sph, sizeof(sph));
+    for (const u64 instruction : lowered.instructions) {
+        if ((code.size() * sizeof(u64)) % 32 == 0) {
+            code.push_back(0); // Maxwell scheduling word, absent from native instructions.
+        }
+        code.push_back(instruction);
+    }
+    // Keep the original padding ABI, then retain the native address map as
+    // unreachable data. Existing disk-cache serialization carries it with code.
+    code.push_back(0);
+    Shader::Ampere::AppendNativeAddressMap(code, lowered.source_instruction_offsets);
+    cached_lowest = 0;
+    cached_highest = static_cast<u32>((code.size() - 1) * sizeof(u64));
+    LOG_INFO(Shader, "Lowered native shader at {:#x}: {} source bytes, {} instructions",
+             program_base, native_source_size, lowered.instructions.size());
+}
+
 GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
                                          Tegra::MemoryManager& gpu_memory_,
                                          Maxwell::ShaderType program, GPUVAddr program_base_,
@@ -298,6 +469,23 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
     : GenericEnvironment{gpu_memory_, program_base_, start_address_}, maxwell3d{&maxwell3d_} {
     gpu_memory->ReadBlock(program_base + start_address, &sph, sizeof(sph));
     initial_offset = sizeof(sph);
+    texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot;
+    if (maxwell3d->GetNativeMacroEngine().native_turing) {
+        if (program != Maxwell::ShaderType::VertexB && program != Maxwell::ShaderType::Pixel) {
+            throw Shader::RuntimeError("Unsupported native graphics shader stage {}", static_cast<u32>(program));
+        }
+        if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) {
+            const u32 native_stage = program == Maxwell::ShaderType::Pixel ? 4 : 0;
+            const auto& table = maxwell3d->state.shader_stages[native_stage].const_buffers[1];
+            if (table.enabled && table.size && table.size <= Tegra::Engines::Maxwell3D::Regs::MaxConstBufferSize) {
+                std::vector<u8> bytes(table.size);
+                gpu_memory->ReadBlock(table.address, bytes.data(), bytes.size());
+                std::ofstream dump(fmt::format("native-graphics-cbuf-{}-1.bin", static_cast<u32>(program)), std::ios::binary);
+                dump.write(reinterpret_cast<const char*>(bytes.data()), bytes.size());
+            }
+        }
+        InitializeNativeCode();
+    }
     gp_passthrough_mask = maxwell3d->regs.post_vtg_shader_attrib_skip_mask;
     switch (program) {
     case Maxwell::ShaderType::VertexA:
@@ -331,8 +519,7 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
     const u64 local_size{sph.LocalMemorySize()};
     ASSERT(local_size <= (std::numeric_limits<u32>::max)());
     local_memory_size = static_cast<u32>(local_size) + sph.common3.shader_local_memory_crs_size;
-    texture_bound = maxwell3d->regs.bindless_texture_const_buffer_slot;
-    is_proprietary_driver = texture_bound == 2;
+    is_proprietary_driver = maxwell3d->GetNativeMacroEngine().native_turing || texture_bound == 2;
     has_hle_engine_state =
         maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro;
 }
@@ -417,6 +604,12 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com
     is_proprietary_driver = texture_bound == 2;
     shared_memory_size = qmd.shared_alloc;
     workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z};
+    if (kepler_compute->IsNativeMode()) {
+        initial_offset = sizeof(sph);
+        is_proprietary_driver = true;
+        InitializeNativeCode(0);
+        start_address = sizeof(sph);
+    }
 }
 
 u32 ComputeEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
@@ -511,7 +704,7 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
     if (stage == Shader::Stage::Compute) {
         file.read(reinterpret_cast<char*>(&workgroup_size), sizeof(workgroup_size))
             .read(reinterpret_cast<char*>(&shared_memory_size), sizeof(shared_memory_size));
-        initial_offset = 0;
+        initial_offset = Shader::Ampere::HasNativeAddressMap(code) ? sizeof(sph) : 0;
     } else {
         file.read(reinterpret_cast<char*>(&sph), sizeof(sph));
         initial_offset = sizeof(sph);
@@ -526,6 +719,13 @@ void FileEnvironment::Dump(u64 pipeline_hash, u64 shader_hash) {
     DumpImpl(pipeline_hash, shader_hash, code, read_highest, read_lowest, initial_offset, stage);
 }
 
+u32 FileEnvironment::TranslateProgramAddress(u32 address) const {
+    if (const auto translated = Shader::Ampere::NativeProgramAddress(code, address)) return *translated;
+    if (Shader::Ampere::HasNativeAddressMap(code))
+        throw Shader::RuntimeError("Invalid cached native indirect branch address {:#x}", address);
+    return address;
+}
+
 u64 FileEnvironment::ReadInstruction(u32 address) {
     if (address < read_lowest || address > read_highest) {
         throw Shader::LogicError("Out of bounds address {}", address);