diff options
Diffstat (limited to 'src')
41 files changed, 2930 insertions, 69 deletions
diff --git a/src/shader_recompiler/CMakeLists.txt b/src/shader_recompiler/CMakeLists.txt index cc9ce95d92..39484a1bd5 100644 --- a/src/shader_recompiler/CMakeLists.txt +++ b/src/shader_recompiler/CMakeLists.txt @@ -58,6 +58,9 @@ add_library(shader_recompiler STATIC frontend/ir/type.h frontend/ir/value.cpp frontend/ir/value.h + frontend/ampere/lower_to_maxwell.h + frontend/ampere/program_override.h + frontend/ampere/program_address_map.h frontend/maxwell/control_flow.cpp frontend/maxwell/control_flow.h frontend/maxwell/decode.cpp diff --git a/src/shader_recompiler/environment.h b/src/shader_recompiler/environment.h index 184b56ebf3..67486430c8 100644 --- a/src/shader_recompiler/environment.h +++ b/src/shader_recompiler/environment.h @@ -74,6 +74,8 @@ public: [[nodiscard]] virtual u64 ReadInstruction(u32 address) = 0; + [[nodiscard]] virtual u32 TranslateProgramAddress(u32 address) const { return address; } + [[nodiscard]] virtual u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) = 0; [[nodiscard]] virtual TextureType ReadTextureType(u32 raw_handle) = 0; diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h new file mode 100644 index 0000000000..524688f036 --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -0,0 +1,1596 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +// Encoding reference: Mesa NAK sm70_encode.rs (MIT), Copyright 2022 Collabora, Ltd. +#pragma once + +#include <algorithm> +#include <array> +#include <bit> +#include <span> +#include <optional> +#include <string_view> +#include <vector> +#include "common/common_types.h" + +namespace Shader::Ampere { + +struct Instruction { + std::array<u32, 4> words; + + u32 Bits(u32 offset, u32 count) const { + u64 bits = words[offset / 32]; + if (offset / 32 + 1 < words.size()) { + bits |= u64{words[offset / 32 + 1]} << 32; + } + return static_cast<u32>((bits >> (offset % 32)) & ((u64{1} << count) - 1)); + } +}; +static_assert(sizeof(Instruction) == 16); + +struct LoweredProgram { + std::vector<u64> instructions; + std::vector<u32> source_instruction_offsets; + u32 source_size{}; + u32 required_texture_buffer = ~0U; + u32 error_offset{}; + std::string_view error; + + explicit operator bool() const { return error.empty(); } +}; + +// Lower only operations whose semantics are representable by the existing +// Maxwell frontend. These are new instructions, not reinterpretations of the +// native bytes. Scheduling/control words are added by the environment adapter. +inline LoweredProgram LowerToMaxwell(std::span<const Instruction> code) { + LoweredProgram result; + struct Relocation { u32 output_index; s64 target; }; + std::vector<Relocation> relocations; + std::vector<u32> barrier_tokens; + std::vector<u32> call_targets; + std::vector<u32> call_links; + bool main_exited = false; + bool uses_uniform_bank = false; + 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)); + } + const auto address_of = [](u32 index) { return 0x50U + index * 8 + ((index + 1) / 3) * 8; }; + const auto fail = [&](u32 index, std::string_view reason) { + result.instructions.clear(); + result.source_instruction_offsets.clear(); + result.error_offset = index * sizeof(Instruction); + result.error = reason; + return result; + }; + const auto finish = [&](u32 index) { + for (const auto& relocation : relocations) { + if (relocation.target > index) return fail(index, "native branch or call leaves captured program"); + const u32 target = result.source_instruction_offsets[static_cast<size_t>(relocation.target)]; + const s64 relative = s64{address_of(target)} - address_of(relocation.output_index) - 8; + if (relative < -0x800000 || relative > 0x7fffff) return fail(index, "lowered branch distance"); + result.instructions[relocation.output_index] |= (static_cast<u64>(relative) & 0xffffff) << 20; + } + result.source_size = (index + 1) * sizeof(Instruction); + return result; + }; + if (code.empty() || code.size() > 65536) return fail(0, "invalid native program size"); + for (u32 index = 0; index < code.size(); ++index) { + auto insn = code[index]; + result.source_instruction_offsets.push_back(static_cast<u32>(result.instructions.size())); + u32 opcode = insn.Bits(0, 12); + std::optional<u32> active_link; + u32 active_entry = 0; + if (main_exited) { + for (size_t call = 0; call < call_targets.size(); ++call) { + if (call_targets[call] <= index && call_targets[call] >= active_entry) { + active_entry = call_targets[call]; + active_link = call_links[call]; + } + } + if (active_link && insn.Bits(16, 8) == *active_link) + return fail(index, "native callee overwrites link register"); + } + // Scratch values and emulated uniform state occupy reserved high GPRs. + // Reject native accesses that could alias that state. + for (u32 bit : {16U, 24U, 32U, 64U}) { + if (bit == 32 && insn.Bits(9, 3) != 1 && opcode != 0xb60 && opcode != 0xb66) continue; + const u32 reg = insn.Bits(bit, 8); + if (reg >= (uses_uniform_bank ? 216U : 246U) && reg < 255) return fail(index, "native program uses lowering scratch registers"); + } + const auto uniform_register = [&](u32 reg) -> u32 { + if (reg < 16) return 216 + reg; + if (reg == 63) return 255; + return ~0U; + }; + if (opcode == 0x2a4 || opcode == 0x891 || opcode == 0x890 || opcode == 0x290 || + opcode == 0x292 || opcode == 0x892 || opcode == 0x291) { + if (insn.Bits(12, 4) != 7) return fail(index, "native uniform arithmetic predication"); + const u32 dst = uniform_register(insn.Bits(16, 8)); + const u32 a = uniform_register(insn.Bits(24, 8)); + const u32 c = uniform_register(insn.Bits(64, 8)); + if (dst == ~0U || a == ~0U || c == ~0U) return fail(index, "native uniform arithmetic register"); + insn.words[0] = (insn.words[0] & 0xffff) | (dst << 16) | (a << 24); + insn.words[2] = (insn.words[2] & ~0x080000ffU) | c; + if (opcode == 0x2a4 || opcode == 0x290 || opcode == 0x292 || opcode == 0x291) { + const u32 b = uniform_register(insn.Bits(32, 8)); + if (b == ~0U) return fail(index, "native uniform arithmetic second register"); + insn.words[1] = (insn.words[1] & ~0xffU) | b; + } + if ((opcode == 0x292 || opcode == 0x892) && insn.Bits(80, 4) != 0xe) + return fail(index, "native uniform logic predicate output"); + opcode = opcode == 0x2a4 ? 0x224 : opcode == 0x891 ? 0x811 : opcode == 0x890 ? 0x810 + : opcode == 0x292 ? 0x212 : opcode == 0x892 ? 0x812 : opcode == 0x291 ? 0x211 : 0x210; + } + const bool uniform_predication = opcode == 0xabb || opcode == 0xab9 || opcode == 0xc82 || opcode == 0x882; + const u32 native_predicate = insn.Bits(12, 4); + const bool guarded_uniform = uniform_predication && native_predicate != 7; + if (guarded_uniform && (native_predicate & 7) != 0) + return fail(index, "native uniform execution predicate"); + 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; }; + const auto cbuf = [&]() -> u64 { + return (u64{insn.Bits(38, 16) / 4} << 20) | (u64{insn.Bits(54, 5)} << 34); + }; + const auto valid_cbuf = [&] { + return !insn.Bits(91, 1) && insn.Bits(54, 5) < 18 && + (insn.Bits(38, 16) % 4) == 0; + }; + std::string_view operand_error; + const auto emit = [&](u64 word) { result.instructions.push_back(word); }; + const auto save_predicate = [&] { + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{0x40} << 20)); + }; + const auto restore_predicate = [&] { + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{0x40} << 20)); + }; + if (guarded_uniform) { + save_predicate(); + emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{0x40} << 20)); + } + const auto mov_imm = [&](u32 reg, u32 value) { + emit(0x0100000000000000ULL | predicate | reg | (u64{15} << 12) | (u64{value} << 20)); + }; + const auto materialize = [&](u32 kind, u32 reg, u32 scratch, bool absolute, bool negate) -> u32 { + if (kind == 3) { + if (reg < 16) reg += 216; + else if (reg == 63) reg = 255; + else { operand_error = "native uniform source register"; return 255; } + } else if (kind == 1) { + mov_imm(scratch, insn.words[1]); + reg = scratch; + } else if (kind == 2) { + if (!valid_cbuf()) { operand_error = "native constant buffer address"; return 255; } + emit(0x4c98000000000000ULL | predicate | scratch | (u64{15} << 39) | cbuf()); + reg = scratch; + } + if (absolute || negate) { + if (reg != scratch) emit(0x5c98000000000000ULL | predicate | scratch | (u64{reg} << 20) | (u64{15} << 39)); + if (absolute) emit(0x0400000000000000ULL | predicate | scratch | (u64{scratch} << 8) | (u64{0x7fffffff} << 20)); + if (negate) emit(0x0400000000000000ULL | predicate | scratch | (u64{scratch} << 8) | (u64{0x80000000} << 20) | (u64{2} << 53)); + reg = scratch; + } + return reg; + }; + const auto operand = [&](u32 which, bool floating = true) -> u32 { + const u32 form = insn.Bits(9, 3); + u32 reg{}, kind{}, abs{}, neg{}; + if (which == 0) { + reg = insn.Bits(24, 8); abs = insn.Bits(73, 1); neg = insn.Bits(72, 1); + } else { + if (form != 1 && form != 2 && form != 3 && form != 4 && form != 5 && form != 6 && form != 7) { + operand_error = "native ALU source form"; return 255; + } + const bool wide = which == 1 ? (form == 4 || form == 5 || form == 6) : (form == 2 || form == 3 || form == 7); + if (wide) { + kind = form == 6 || form == 7 ? 3 : (form == 4 || form == 2) ? 1 : 2; + if (kind == 3) reg = insn.Bits(32, 8); + if (kind == 2 || kind == 3) { abs = insn.Bits(62, 1); neg = insn.Bits(63, 1); } + } else { + const bool upper = which == 1 ? (form == 2 || form == 3 || form == 7) : true; + reg = insn.Bits(upper ? 64 : 32, 8); + abs = insn.Bits(upper ? 74 : 62, 1); neg = insn.Bits(upper ? 75 : 63, 1); + } + } + return materialize(kind, reg, 250 + which, floating && abs, floating && neg); + }; + struct HalfSource { u32 reg, swizzle; bool absolute, negate; }; + const auto half_source = [&](u32 which) -> HalfSource { + if (which == 0) return {insn.Bits(24, 8), insn.Bits(74, 2), + insn.Bits(73, 1) != 0, insn.Bits(72, 1) != 0}; + const u32 form = insn.Bits(9, 3); + const bool wide = which == 1 ? (form == 4 || form == 5) : (form == 2 || form == 3); + if (wide) { + const bool immediate = form == 2 || form == 4; + const u32 reg = materialize(immediate ? 1 : 2, 255, 250 + which, false, false); + return {reg, immediate ? 0U : insn.Bits(60, 2), + !immediate && insn.Bits(62, 1), !immediate && insn.Bits(63, 1)}; + } + const bool upper = which == 2 || form == 2 || form == 3; + return {insn.Bits(upper ? 64 : 32, 8), insn.Bits(upper ? 81 : 60, 2), + insn.Bits(upper ? 83 : 62, 1) != 0, insn.Bits(upper ? 84 : 63, 1) != 0}; + }; + const auto extract_half = [&](HalfSource source, u32 lane, u32 target) { + const bool high = source.swizzle == 3 || (source.swizzle == 0 && lane == 1); + emit(0x5ca8000000000000ULL | predicate | target | (u64{2} << 8) | + (u64{source.swizzle == 1 ? 2U : 1U} << 10) | (u64{source.reg} << 20) | + (u64{high} << 41) | (u64{source.negate} << 45) | (u64{source.absolute} << 49)); + }; + const auto xmad = [&](u32 target, u32 a, u32 b, u32 c, bool a_high, + bool b_high, u32 mode, bool immediate) { + emit((immediate ? 0x3600000000000000ULL : 0x5b00000000000000ULL) | + predicate | target | (u64{a} << 8) | + (u64{immediate ? ((b >> (b_high ? 16 : 0)) & 0xffff) : b} << 20) | + (u64{!immediate && b_high} << 35) | (u64{c} << 39) | + (u64{mode} << 50) | (u64{a_high} << 53)); + }; + const auto multiply_low = [&](u32 a, u32 b, u32 c, bool immediate) { + xmad(250, a, b, c, false, false, 0, immediate); + xmad(251, a, b, 255, true, false, 0, immediate); + xmad(251, a, b, 251, false, true, 0, immediate); + emit(base(0x5c18000000000000ULL) | (u64{251} << 8) | + (u64{250} << 20) | (u64{16} << 39)); + }; + switch (opcode) { + case 0x304: { // F2F.F16.F32, scalar round-to-nearest conversion. + if (insn.words[2] != 0x00200800 || insn.Bits(24, 8) || + (insn.words[1] & ~0xc00000ffU)) + return fail(index, "native scalar float conversion mode"); + emit(base(0x5ca8000000000000ULL) | (u64{1} << 8) | (u64{2} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(62, 1)} << 49) | + (u64{insn.Bits(63, 1)} << 45)); + break; + } + case 0x23e: case 0xa3e: { // F2FP.PACK_AB / MERGE_C, round-to-nearest + if (insn.Bits(62, 2) || insn.Bits(73, 5) || insn.Bits(79, 2) || + insn.Bits(81, 10) || insn.Bits(91, 1)) { + return fail(index, "native float packing modifiers"); + } + const bool merge = insn.Bits(78, 1); + if (insn.Bits(72, 1)) return fail(index, "native float packing selector"); + const u32 a = merge ? operand(2, false) : operand(0, false); + const u32 b = operand(1, false); + // Canonical F2F writes the converted half in its low 16 bits. + if (merge) { + emit(0x5c98000000000000ULL | predicate | 250 | (u64{a} << 20) | (u64{15} << 39)); + } else { + emit(0x5ca8000000000000ULL | predicate | 250 | (u64{1} << 8) | + (u64{2} << 10) | (u64{a} << 20)); + } + emit(0x5ca8000000000000ULL | predicate | 251 | (u64{1} << 8) | + (u64{2} << 10) | (u64{b} << 20)); + emit(base(0x5c18000000000000ULL) | (u64{250} << 8) | (u64{251} << 20) | (u64{16} << 39)); + break; + } + case 0x232: case 0x832: case 0xa32: { // HMUL2 + if (insn.Bits(78, 2) || insn.Bits(81, 10) || insn.Bits(91, 1)) + return fail(index, "native half multiply mode"); + const auto a = half_source(0), b = half_source(1); + const u32 precision = insn.Bits(76, 1) ? 2 : insn.Bits(80, 1); + emit(base(0x5d08000000000000ULL) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{b.absolute} << 30) | (u64{a.negate != b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{precision} << 39) | + (u64{a.absolute} << 44) | (u64{a.swizzle} << 47)); + break; + } + case 0x230: case 0x430: case 0x630: { // HADD2 + if (insn.Bits(76, 1) || insn.Bits(79, 1) || insn.Bits(81, 10) || insn.Bits(91, 1)) + return fail(index, "native half add mode"); + const auto a = half_source(0); + const auto b = half_source(opcode == 0x230 ? 1 : 2); + if (insn.Bits(78, 1)) { + if (a.reg != 255 || !a.negate || a.absolute || a.swizzle != 0 || + b.swizzle < 2 || insn.Bits(77, 1) || insn.Bits(80, 1)) + return fail(index, "native float-output half add mode"); + extract_half(b, 0, static_cast<u32>(dest)); + } else { + emit(base(0x5d10000000000000ULL) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{b.absolute} << 30) | (u64{b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{insn.Bits(80, 1)} << 39) | + (u64{a.negate} << 43) | (u64{a.absolute} << 44) | (u64{a.swizzle} << 47)); + } + break; + } + case 0x231: case 0x431: case 0x631: case 0x831: case 0xa31: { // HFMA2 + const bool relu = insn.Bits(79, 1) != 0; + if (insn.Bits(78, 1) || insn.Bits(85, 1) || insn.Bits(91, 1) || + (relu ? (insn.Bits(87, 4) != 7 || + insn.Bits(77, 1)) : insn.Bits(87, 4) != 0)) + return fail(index, "native half fused multiply-add mode"); + if (relu && native_predicate != 7) { + // Guard the whole expansion before its temporary P6 updates. + relocations.push_back({static_cast<u32>(result.instructions.size()), index + 1}); + emit(0xe24000000000000fULL | (u64{native_predicate ^ 8U} << 16)); + predicate = u64{7} << 16; + } + auto a = half_source(0); + auto b = half_source(1); + const auto c = half_source(2); + if (insn.Bits(86, 1)) { // H0_NH1: negate only the upper packed component. + if (b.swizzle != 0 || b.absolute || b.reg >= 250) + return fail(index, "native half mixed-sign source mode"); + emit(0x0400000000000000ULL | predicate | 251 | (u64{b.reg} << 8) | + (u64{0x80000000} << 20) | (u64{2} << 53)); + b.reg = 251; + } + if (b.absolute || c.absolute || (a.absolute && a.reg >= 248)) + return fail(index, "native half fused multiply-add absolute source"); + if (a.absolute) { + emit(0x0400000000000000ULL | predicate | 250 | (u64{a.reg} << 8) | + (u64{a.swizzle == 1 ? 0x7fffffffU : 0x7fff7fffU} << 20)); + a.reg = 250; + a.absolute = false; + } + const u32 precision = insn.Bits(76, 1) ? 2 : insn.Bits(80, 1); + emit(0x5d00000000000000ULL | predicate | (relu ? 254U : dest) | (u64{a.reg} << 8) | (u64{b.reg} << 20) | + (u64{b.swizzle} << 28) | (u64{c.negate} << 30) | (u64{a.negate != b.negate} << 31) | + (u64{insn.Bits(77, 1)} << 32) | (u64{c.swizzle} << 35) | + (u64{precision} << 37) | (u64{c.reg} << 39) | (u64{a.swizzle} << 47)); + if (relu) { + // Clamp the rounded binary16 lanes, preserving positive infinity + // and canonicalizing either sign of NaN without saturating to one. + save_predicate(); + for (u32 lane = 0; lane < 2; ++lane) { + emit(0x3828000000070000ULL | 251 | (u64{254} << 8) | (u64{lane * 16} << 20)); + emit(0x0400000000070000ULL | 251 | (u64{251} << 8) | (u64{0xffff} << 20)); + emit(0x0400000000070000ULL | 252 | (u64{251} << 8) | (u64{0x7fff} << 20)); + mov_imm(250, 0x7c00); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39) | (u64{1} << 49)); + emit(0x3828000000070000ULL | 252 | (u64{251} << 8) | (u64{15} << 20)); + mov_imm(250, 1); + emit(0x5c10000000070000ULL | 252 | (u64{252} << 8) | (u64{250} << 20) | (u64{1} << 48)); + emit(0x5c40000000070000ULL | 251 | (u64{251} << 8) | (u64{252} << 20) | (u64{7} << 48)); + mov_imm(250, 0x7fff); + emit(0x5ca0000000070000ULL | 251 | (u64{250} << 8) | (u64{251} << 20) | (u64{6} << 39)); + if (lane == 0) { + emit(base(0x5c98000000000000ULL) | (u64{251} << 20) | (u64{15} << 39)); + } else { + emit(0x3848000000070000ULL | 251 | (u64{251} << 8) | (u64{16} << 20)); + emit(base(0x5c40000000000000ULL) | (dest << 8) | (u64{251} << 20) | + (u64{1} << 41) | (u64{7} << 48)); + } + } + restore_predicate(); + } + break; + } + case 0x240: case 0x840: { // HMNMX2, exact half promotion for min/max + if (insn.Bits(76, 11) || insn.Bits(91, 1)) + return fail(index, "native half min/max mode"); + const auto a = half_source(0); + const auto b = opcode == 0x840 ? HalfSource{255, 0, false, false} : half_source(1); + if (a.swizzle == 1 || b.swizzle == 1 || b.reg >= 250 && b.reg != 255) + return fail(index, "native half min/max source mode"); + for (u32 lane = 0; lane < 2; ++lane) { + extract_half(a, lane, 250); + if (opcode == 0x840) { + const u32 half = (insn.words[1] >> (lane * 16)) & 0xffff; + mov_imm(251, half); + extract_half(HalfSource{251, 2, false, false}, 0, 251); + } else { + extract_half(b, lane, 251); + } + const u32 target = lane == 0 ? 252 : 250; + emit(0x5c60000000000000ULL | predicate | target | (u64{250} << 8) | (u64{251} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(0x5ca8000000000000ULL | predicate | target | (u64{1} << 8) | + (u64{2} << 10) | (u64{target} << 20)); + } + emit(base(0x5c18000000000000ULL) | (u64{250} << 8) | (u64{252} << 20) | (u64{16} << 39)); + break; + } + case 0xc24: case 0xe24: { // IMAD low word with a uniform multiplier or addend. + if ((insn.words[2] & ~0x080002ffU) != 0x078e0000 || (insn.words[1] & ~0xffU)) + return fail(index, "native uniform-addend multiply mode"); + const u32 a = operand(0, false), b = operand(1, false), c = operand(2, false); + multiply_low(a, b, c, false); + break; + } + case 0x224: { // IMAD register / zero-product move + if (insn.Bits(24, 8) == 255 && insn.words[1] == 255 && insn.Bits(75, 1) && + (insn.words[2] & ~0xaffU) == 0x078e0000) { + emit(base(0x5c10000000000000ULL) | (u64{255} << 8) | + (u64{insn.Bits(64, 8)} << 20) | (u64{1} << 48)); + break; + } + if (insn.Bits(24, 8) == 255 && insn.words[1] == 255 && + (insn.words[2] & ~0xffU) == 0x078e0000) { + emit(base(0x5c98000000000000ULL) | (u64{15} << 39) | (u64{insn.Bits(64, 8)} << 20)); + break; + } + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || (insn.words[1] & ~0xffU)) + return fail(index, "native integer register multiply mode"); + multiply_low(insn.Bits(24, 8), insn.Bits(32, 8), insn.Bits(64, 8), false); + break; + } + case 0x827: { // IMAD.HI.U32, captured zero addend + if (insn.words[2] != 0x078e00ff) return fail(index, "native high multiply mode"); + const u32 a = insn.Bits(24, 8), b = insn.words[1]; + // Base-2^16 long multiplication. Each intermediate sum fits U32. + xmad(250, a, b, 255, false, false, 0, true); + xmad(251, a, b, 250, true, false, 2, true); + xmad(250, a, b, 251, false, true, 1, true); + xmad(252, a, b, 251, true, true, 2, true); + emit(0x3828000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{16} << 20)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{252} << 20)); + break; + } + case 0xa24: { // IMAD, constant-buffer multiplier + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || !valid_cbuf() || insn.Bits(62, 2)) + return fail(index, "native integer constant-buffer multiply mode"); + emit(0x4c98000000000000ULL | predicate | 252 | (u64{15} << 39) | cbuf()); + multiply_low(insn.Bits(24, 8), 252, insn.Bits(64, 8), false); + break; + } + case 0x424: case 0x624: { // IMAD.MOV.U32: zero product plus immediate/cbuf + if (opcode == 0x624 && insn.Bits(24, 8) == 255 && + (insn.words[2] & ~0x07800000U) == 0x000e06ffU) { + const u32 source = operand(2, false); + mov_imm(250, 1); + emit(0x5ca0000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(base(0x5c10000000000000ULL) | (u64{source} << 8) | (u64{250} << 20)); + break; + } + if (insn.Bits(24, 8) != 255 || insn.Bits(64, 8) != 255 || insn.words[2] != 0x078e00ff) { + if ((insn.words[2] & ~0x2ffU) != 0x078e0000 || (opcode == 0x624 && insn.Bits(62, 2))) + return fail(index, "native multiply wide addend mode"); + const u32 a = operand(0, false), b = operand(1, false), c = operand(2, false); + multiply_low(a, b, c, false); + break; + } + if (insn.Bits(24, 8) != 255 || insn.Bits(64, 8) != 255 || + insn.words[2] != 0x078e00ff) + return fail(index, "native integer move mode"); + if (opcode == 0x624) { + if (!valid_cbuf() || insn.Bits(62, 2)) + return fail(index, "native integer move constant buffer"); + emit(base(0x4c98000000000000ULL) | (u64{15} << 39) | cbuf()); + } else { + emit(base(0x0100000000000000ULL) | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); + } + break; + } + case 0x824: { // Captured low-word IMAD forms + if (insn.words[1] == 1 && (insn.words[2] & 0xffffff00U) == 0x078e0200) { + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(64, 8)} << 20)); + break; + } + if (insn.words[1] != 48 || insn.words[2] != 0x078e02ff) { + if ((insn.words[2] & ~0x2ffU) != 0x078e0000) + return fail(index, "native integer multiply mode"); + multiply_low(insn.Bits(24, 8), insn.words[1], insn.Bits(64, 8), true); + break; + } + const u64 source = insn.Bits(24, 8); + emit(0x5c18000000000000ULL | predicate | 250 | (source << 8) | (source << 20) | (u64{1} << 39)); + emit(base(0x3848000000000000ULL) | (u64{250} << 8) | (u64{4} << 20)); + break; + } + case 0x806: { // VOTE: register ballot, ordinary predicate input and output. + if (insn.words[1] || insn.Bits(72, 2) > 2 || (insn.words[2] & ~0x078e0300U)) + return fail(index, "native vote mode"); + emit(base(0x50d8000000000000ULL) | (u64{insn.Bits(87, 4)} << 39) | + (u64{insn.Bits(81, 3)} << 45) | (u64{insn.Bits(72, 2)} << 48)); + break; + } + case 0x886: { // VOTEU ballot and independent uniform predicate output. + const u32 output = insn.Bits(81, 3), target = insn.Bits(16, 8); + if ((target >= 16 && target != 63) || insn.words[1] || + (output > 1 && output != 7) || insn.Bits(72, 2) > 2 || native_predicate != 7) + return fail(index, "native uniform vote mode"); + const u32 temporary = output == 7 ? 7 : 6 - output; + const u32 mask = 1U << temporary; + if (output != 7) emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x50d8000000070000ULL | uniform_register(target) | + (u64{insn.Bits(87, 4)} << 39) | (u64{temporary} << 45) | (u64{insn.Bits(72, 2)} << 48)); + if (output != 7) { + emit(0x38e8000000070000ULL | 239 | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + } + break; + } + case 0x3c2: { // R2UR: broadcast the first lane satisfying the execution predicate. + if (dest >= 16 || insn.words[1] || insn.words[2] != 0x000e0000) + return fail(index, "native register-to-uniform mode"); + save_predicate(); + emit(0x50d8000000070000ULL | 250 | (u64{native_predicate} << 39) | + (u64{6} << 45) | (u64{1} << 48)); + emit(0x5c10000000070000ULL | 251 | (u64{250} << 8) | (u64{255} << 20) | (u64{1} << 49)); + emit(0x5c40000000070000ULL | 250 | (u64{250} << 8) | (u64{251} << 20) | (u64{7} << 48)); + emit(0x5c30000000070000ULL | 250 | (u64{250} << 20)); + emit(0xef10000000070000ULL | 251 | (u64{insn.Bits(24, 8)} << 8) | + (u64{250} << 20) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5ca0000000070000ULL | (dest + 216) | (u64{251} << 8) | + (u64{dest + 216} << 20) | (u64{6} << 39)); + restore_predicate(); + break; + } + case 0x81c: { // Observed predicate LUT constants and UP0-to-predicate copy. + const u32 target = insn.Bits(81, 3); + const u32 lut = insn.Bits(64, 3) | (insn.Bits(72, 5) << 3); + if ((native_predicate != 7 && insn.Bits(67, 1)) || target == 7 || insn.Bits(84, 3) != 7 || + insn.Bits(16, 8) || insn.words[1]) + return fail(index, "native predicate LUT mode"); + if (!insn.Bits(67, 1) && (lut == 0xfe || lut == 0xa8)) { + emit(0x5090000000000000ULL | predicate | (u64{target} << 3) | 7 | + (u64{insn.Bits(87, 4)} << 12) | (u64{1} << 24) | + (u64{insn.Bits(77, 4)} << 29) | (u64{insn.Bits(68, 4)} << 39) | + (u64{lut == 0xfe} << 45)); + break; + } + if (!insn.Bits(67, 1) && (lut == 0x80 || + (lut == 8 && insn.Bits(68, 4) == 7 && insn.Bits(77, 4) == 7))) { + emit(0x5090000000000000ULL | predicate | (u64{target} << 3) | 7 | + (u64{insn.Bits(87, 4) ^ (lut == 8 ? 8U : 0U)} << 12) | (u64{insn.Bits(77, 4)} << 29) | + (u64{insn.Bits(68, 4)} << 39)); + break; + } + if (insn.Bits(77, 4) != 7 || insn.Bits(87, 4) != 7) + return fail(index, "native predicate LUT source mode"); + if (insn.Bits(67, 1) && insn.Bits(68, 4) == 0 && lut == 0x80) { + emit(0x3828000000070000ULL | 250 | (u64{239} << 8) | (u64{6} << 20)); + emit(0x3848000000070000ULL | 250 | (u64{250} << 8) | (u64{target} << 20)); + } else if (!insn.Bits(67, 1) && insn.Bits(68, 4) == 7 && (lut == 0x80 || lut == 8)) { + emit(0x010000000007f000ULL | 250 | ((lut == 0x80 ? u64{1} << target : 0ULL) << 20)); + } else return fail(index, "native predicate LUT truth table"); + emit(0x38f0000000070000ULL | (u64{250} << 8) | (u64{1} << (target + 20))); + break; + } + case 0x882: { // UMOV immediate. + if (dest >= 16 || insn.words[2]) return fail(index, "native uniform immediate move mode"); + mov_imm(static_cast<u32>(dest) + 216, insn.words[1]); + break; + } + case 0xc82: { // UMOV UR, UR (or URZ), including uniform execution predicates. + const u32 source = insn.words[1] < 64 ? uniform_register(insn.words[1]) : ~0U; + if (dest >= 16 || source == ~0U || insn.words[2] != 0x08000000) + return fail(index, "native uniform move mode"); + emit(0x5c98000000000000ULL | predicate | (dest + 216) | + (u64{source} << 20) | (u64{15} << 39)); + break; + } + case 0xab9: case 0xabb: { // ULDC U32/U64, direct or indexed. + const u32 uniform_dest = insn.Bits(16, 8), index_reg = insn.Bits(24, 8); + 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) || + (insn.words[2] & ~0xe00U) != (opcode == 0xab9 ? 0U : 0x08000000U) || + (opcode == 0xabb && index_reg >= 16)) + return fail(index, "native uniform constant load mode"); + if (opcode == 0xab9 && width == 4) { + emit(0x4c98000000000000ULL | predicate | (uniform_dest + 216) | (u64{15} << 39) | cbuf()); + } 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)); + } + break; + } + case 0x899: { // USHF.L.U32; accept a zero word outside the selected half. + const u32 dst = insn.Bits(16, 8), low = insn.Bits(24, 8), high = insn.Bits(64, 8); + if (insn.words[2] == (0x08011400U | high) && low == 63 && high < 16 && + dst < 16 && native_predicate == 7 && insn.words[1] <= 32) { + emit(0x3828000000000000ULL | predicate | (dst + 216) | + (u64{high + 216} << 8) | (u64{insn.words[1]} << 20) | (u64{1} << 48)); + break; + } + const bool high_result = insn.words[2] == (0x08010600U | high) && low == 63 && high < 16; + const bool low_result = insn.words[2] == 0x0800063f && low < 16; + if (dst >= 16 || native_predicate != 7 || insn.words[1] > 32 || (!high_result && !low_result)) + return fail(index, "native uniform shift mode"); + emit(0x3848000000000000ULL | predicate | (dst + 216) | + (u64{(high_result ? high : low) + 216} << 8) | (u64{insn.words[1]} << 20)); + break; + } + 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 || + (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)); + break; + } + case 0xf89: { // SHFL with immediate lane and segmentation/clamp mask, any mode. + const u32 lane = insn.Bits(53, 5), mask = insn.Bits(40, 13), mode = insn.Bits(58, 2); + if (insn.words[2] != 0x000e0000 || insn.Bits(32, 8) || insn.Bits(60, 4)) + return fail(index, "native immediate shuffle mode"); + emit(base(0xef10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{lane} << 20) | (u64{1} << 28) | (u64{1} << 29) | + (u64{mode} << 30) | (u64{mask} << 34) | (u64{7} << 48)); + break; + } + case 0x822: { // FSWZADD, bounded quad derivative sign patterns. + const u32 swizzle = insn.words[1]; + if ((swizzle != 0x66 && swizzle != 0x5a) || insn.words[2] > 255) + return fail(index, "native swizzled add mode"); + // Maxwell encodes operand signs rather than the native operation enum. + emit(base(0x50f8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(64, 8)} << 20) | (u64{swizzle == 0x66 ? 0x99U : 0xa5U} << 28)); + break; + } + case 0xb6c: { // TXD ARRAY_2D, four channels and consecutive split destinations. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 coordinates = insn.Bits(24, 8), derivatives = insn.Bits(32, 8); + if ((insn.words[2] & ~0xffU) != 0x001e0f00 || insn.Bits(59, 2) || + insn.Bits(61, 3) != 5 || binding >= 18 || offset > 0x1fff || + (dest & 1) || dest >= 248 || insn.Bits(64, 8) != dest + 2 || + coordinates >= 248 || derivatives >= 247) + return fail(index, "native array gradient texture mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native texture descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(base(0xde00000000000000ULL) | (u64{coordinates} << 8) | + (u64{derivatives} << 20) | (u64{3} << 28) | (u64{15} << 31) | + (u64{offset} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51)); + 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"); + emit(0xef10000000000000ULL | predicate | dest | + (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(32, 8)} << 20) | + (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)) + return fail(index, "native global reduction mode"); + emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | + (u64{insn.Bits(24, 8)} << 8) | (u64{1} << 48)); + break; + } + case 0x2bd: // UFLO.U32, the highest set-bit index (.SH, bit 74: its shift amount). + if (dest >= 16 || insn.words[1] >= 16 || (insn.words[2] & ~0x400U) != 0x080e0000 || native_predicate != 7) + return fail(index, "native uniform leading bit mode"); + emit(0x5c30000000000000ULL | predicate | (dest + 216) | (u64{insn.words[1] + 216} << 20) | + (u64{insn.Bits(74, 1)} << 41)); + break; + case 0x2be: // UBREV, as a bit-reversing zero shift like BREV. + if (dest >= 16 || insn.words[1] >= 16 || insn.words[2] != 0x08000000 || native_predicate != 7) + return fail(index, "native uniform bit reverse mode"); + emit(0x3828000000000000ULL | predicate | (dest + 216) | (u64{insn.words[1] + 216} << 8) | + (u64{1} << 39) | (u64{1} << 40)); + break; + case 0x3a1: { // MATCH.ALL.U64 with only the all-equal predicate retained. + const u32 source = insn.Bits(24, 8), target = insn.Bits(81, 3); + if (dest != 255 || source >= 254 || (source & 1) || target >= 6 || + native_predicate != 7 || insn.words[1] || (insn.words[2] & ~0xe0000U) != 0x200) + return fail(index, "native warp match mode"); + save_predicate(); + emit(0x50d8000000070000ULL | 250 | (u64{7} << 39) | (u64{7} << 45) | (u64{1} << 48)); + emit(0x5c10000000070000ULL | 251 | (u64{250} << 8) | (u64{255} << 20) | (u64{1} << 49)); + emit(0x5c40000000070000ULL | 250 | (u64{250} << 8) | (u64{251} << 20) | (u64{7} << 48)); + emit(0x5c30000000070000ULL | 250 | (u64{250} << 20)); + for (u32 word = 0; word < 2; ++word) + emit(0xef10000000070000ULL | (251 + word) | (u64{source + word} << 8) | + (u64{250} << 20) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{source} << 8) | + (u64{251} << 20) | (u64{7} << 39) | (u64{2} << 49)); + emit(0x5b60000000070000ULL | 7 | (u64{6} << 3) | (u64{source + 1} << 8) | + (u64{252} << 20) | (u64{6} << 39) | (u64{2} << 49)); + emit(0x50d8000000070000ULL | 255 | (u64{6} << 39) | (u64{target} << 45)); + restore_predicate(); + break; + } + case 0x3c4: { // REDUX.SUM over the actual active lanes, modulo 2^32. + const u32 source = insn.Bits(24, 8); + if (dest >= 16 || insn.words[1] || insn.words[2] != 0xc000 || native_predicate != 7) + return fail(index, "native warp sum mode"); + emit(0x50d8000000070000ULL | 250 | (u64{7} << 39) | (u64{7} << 45) | (u64{1} << 48)); + mov_imm(static_cast<u32>(dest) + 216, 0); + for (u32 lane = 0; lane < 32; ++lane) { + emit(0xef10000000070000ULL | 251 | (u64{source} << 8) | (u64{lane} << 20) | + (u64{1} << 28) | (u64{1} << 29) | (u64{31} << 34) | (u64{7} << 48)); + emit(0x3828000000070000ULL | 252 | (u64{250} << 8) | (u64{lane} << 20)); + emit(0x0400000000070000ULL | 252 | (u64{252} << 8) | (u64{1} << 20)); + emit(0x5c10000000070000ULL | 252 | (u64{255} << 8) | (u64{252} << 20) | (u64{1} << 48)); + emit(0x5c40000000070000ULL | 251 | (u64{251} << 8) | (u64{252} << 20) | (u64{7} << 48)); + emit(0x5c10000000070000ULL | (dest + 216) | (u64{dest + 216} << 8) | (u64{251} << 20)); + } + break; + } + case 0xf8c: { // ATOMS.POPC.INC.32 with no returned value: one increment per active lane. + const u32 stride = insn.Bits(78, 2); + if (dest != 255 || insn.words[1] || (insn.words[2] & ~0xc000U) != 0x0d00003f) + return fail(index, "native shared increment mode"); + u32 address = insn.Bits(24, 8); + if (stride && address != 255) { + emit(0x3848000000000000ULL | predicate | 251 | (u64{address} << 8) | + (u64{stride + 1} << 20)); + address = 251; + } + mov_imm(250, 1); + emit(0xec00000000000000ULL | predicate | 255 | (u64{address} << 8) | (u64{250} << 20)); + break; + } + case 0x38c: { // ATOMS.ADD.U32, naturally aligned shared byte offset. + const u32 offset = insn.Bits(40, 24); + if (insn.words[2] || (offset & 3)) return fail(index, "native shared atomic mode"); + emit(0xec00000000000000ULL | predicate | dest | + (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(32, 8)} << 20) | + (u64{offset >> 2} << 30)); + break; + } + case 0x9a8: { // ATOMG.E.ADD.STRONG.GPU U32 through a uniform pointer. + const u32 address = uniform_register(insn.Bits(64, 6)); + if (insn.Bits(24, 8) != 255 || insn.Bits(40, 24) || + (insn.words[2] & ~0x3fU) != 0x081ee140 || address == ~0U || (address & 1)) + return fail(index, "native global atomic mode"); + emit(0xed00000000000000ULL | predicate | dest | (u64{address} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{1} << 48)); + break; + } + case 0xb97: { // SULD.P.2D.IGN, direct descriptor, channel mask at 72-75. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 coordinate = insn.Bits(24, 8); + const u32 channels = insn.Bits(72, 4); + if ((insn.words[2] & ~0xf00U) != 0x001e0000 || channels == 0 || insn.Bits(61, 3) != 3 || + binding >= 18 || offset > 0x1fff || insn.Bits(59, 2) || dest == 255 || coordinate > 253) + return fail(index, "native formatted surface load mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native surface descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(base(0xeb00000000000000ULL) | (u64{coordinate} << 8) | + (u64{channels} << 20) | (u64{3} << 33) | (u64{offset} << 36) | (u64{1} << 51)); + break; + } + case 0xb9b: { // SUST.P.2D/3D.IGN, direct descriptor and all four channels. + const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); + const u32 data = insn.Bits(32, 8), coordinate = insn.Bits(24, 8); + const u32 dimension = insn.Bits(61, 3); + if (insn.words[2] != 0x00100f00 || (dimension != 3 && dimension != 5) || + binding >= 18 || offset > 0x1fff || insn.Bits(59, 2) || data > 252 || coordinate > (dimension == 5 ? 252U : 253U)) + return fail(index, "native formatted surface store mode"); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native surface descriptor buffer conflict"); + result.required_texture_buffer = binding; + emit(0xeb20000000000000ULL | predicate | data | (u64{coordinate} << 8) | + (u64{15} << 20) | (u64{dimension} << 33) | (u64{offset} << 36) | (u64{1} << 51)); + break; + } + case 0x301: // BREV register. + if (insn.words[2] || insn.words[1] > 255) return fail(index, "native bit reverse mode"); + emit(base(0x3828000000000000ULL) | (u64{insn.words[1]} << 8) | + (u64{1} << 39) | (u64{1} << 40)); + break; + case 0x219: { // U32 register shifts with a zero word outside the selected result. + const u32 low = insn.Bits(24, 8), high = insn.Bits(64, 8), shift = insn.Bits(32, 8); + const bool right = insn.Bits(76, 1) != 0, high_result = insn.Bits(80, 1) != 0; + if (insn.words[1] > 255 || insn.Bits(73, 2) != 3 || + (insn.words[2] & ~0x000118ffU) != 0x00000600U || + (right ? (!high_result || low != 255) : (high_result || high != 255))) + return fail(index, "native register shift mode"); + emit(base(right ? 0x5cf8000000000000ULL : 0x5bf8000000000000ULL) | + (u64{right ? high : 255U} << 8) | (u64{shift} << 20) | + (u64{right ? 255U : low} << 39) | (u64{insn.Bits(75, 1)} << 50)); + break; + } + case 0x300: case 0xd00: { // FLO, including a uniform source. + if ((insn.words[2] & ~(0x600U | (opcode == 0xd00 ? 0x08000000U : 0U))) != 0x000e0000U) + return fail(index, "native leading bit mode"); + const u32 source = operand(1, false); + emit(base(0x5c30000000000000ULL) | (u64{source} << 20) | + (u64{insn.Bits(63, 1)} << 40) | (u64{insn.Bits(74, 1)} << 41) | + (u64{insn.Bits(73, 1)} << 48)); + break; + } + case 0x309: case 0xd09: { // POPC, including a uniform source. + if (insn.words[2] != (opcode == 0xd09 ? 0x08000000U : 0U)) + return fail(index, "native population count mode"); + const u32 source = operand(1, false); + emit(base(0x5c08000000000000ULL) | (u64{source} << 20) | (u64{insn.Bits(63, 1)} << 40)); + break; + } + case 0x918: + if (insn.words[1] || insn.words[2]) return fail(index, "native NOP mode"); + emit(0x50b0000000000f00ULL | predicate); + break; + case 0x992: { // MEMBAR.SC: preserve workgroup/device visibility. + const u32 scope = insn.Bits(76, 3); + if (insn.words[1] || insn.words[2] != (scope << 12) || + (scope != 0 && scope != 2 && scope != 3 && scope != 5)) + return fail(index, "native memory barrier scope"); + const u32 legacy_scope = scope == 0 ? 0 : scope == 2 ? 1 : scope == 3 ? 2 : 3; + emit(0xef98000000000000ULL | predicate | (u64{legacy_scope} << 8)); + break; + } + case 0x9ab: // ERRBAR: wait for prior memory operations on the host backend. + if (insn.words[1] || insn.words[2]) return fail(index, "native error barrier mode"); + emit(0xef98000000000000ULL | predicate | (u64{1} << 8)); + break; + case 0xb1d: // BAR.SYNC 0, full workgroup. + if (insn.words[1] || insn.words[2] || native_predicate != 7) + return fail(index, "native workgroup barrier mode"); + emit(0xf0a8000000000000ULL | predicate | (u64{1} << 43) | + (u64{1} << 44) | (u64{1} << 39) | (u64{7} << 39)); + break; + case 0x388: case 0x984: { // Shared memory load/store with scaled register offsets. + const u32 width = insn.Bits(73, 3), stride = insn.Bits(78, 2); + if (width > 6 || insn.words[2] != ((width << 9) | (stride << 14)) || + (opcode == 0x984 && insn.Bits(32, 8))) + return fail(index, "native shared memory addressing mode"); + u32 address = insn.Bits(24, 8); + if (stride && address != 255) { + emit(0x3848000000000000ULL | predicate | 250 | + (u64{address} << 8) | (u64{stride + 1} << 20)); + address = 250; + } + const u32 data = opcode == 0x388 ? insn.Bits(32, 8) : static_cast<u32>(dest); + emit((opcode == 0x388 ? 0xef58000000000000ULL : 0xef48000000000000ULL) | + predicate | data | (u64{address} << 8) | (u64{insn.Bits(40, 24)} << 20) | + (u64{width} << 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) || + (special != 0 && special != 33 && special != 34 && special != 35 && + special != 37 && special != 38 && special != 39 && + special != 56 && special != 57 && special != 58 && special != 59 && special != 60)) + return fail(index, "native special register"); + const u32 target = opcode == 0x9c3 ? uniform_register(static_cast<u32>(dest)) : static_cast<u32>(dest); + if (target == ~0U) return fail(index, "native uniform special destination"); + emit(0xf0c8000000000000ULL | predicate | target | (u64{special} << 20)); + break; + } + case 0x28c: case 0x88c: { // UISETP, signed (bit 73) or unsigned, with UPT accumulation. + const u32 output = insn.Bits(81, 3); + const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x88c ? 252U : uniform_register(insn.Bits(32, 8)); + if (output > 1 || insn.Bits(84, 3) != 7 || insn.Bits(87, 4) != 7 || + (insn.Bits(72, 4) & ~2U) || insn.Bits(12, 4) != 7 || a == ~0U || b == ~0U) + return fail(index, "native uniform comparison mode"); + const u32 temporary = 6 - output, mask = 1U << temporary; + if (opcode == 0x88c) mov_imm(252, insn.words[1]); + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x5b60000000070000ULL | 7 | (u64{temporary} << 3) | + (u64{a} << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{insn.Bits(73, 1)} << 48) | + (u64{insn.Bits(76, 3)} << 49)); + emit(0x38e8000000070000ULL | 239 | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + break; + } + case 0x287: case 0x887: { // USEL, UP0/UP1; preserve ordinary predicates. + const u32 condition = insn.Bits(87, 3); + const u32 target = uniform_register(static_cast<u32>(dest)); + const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x887 ? 252U : uniform_register(insn.Bits(32, 8)); + if (condition > 1 || target == ~0U || a == ~0U || b == ~0U || + insn.Bits(12, 4) != 7 || (insn.words[2] & ~0x07800000U) != 0x08000000U) + return fail(index, "native uniform select mode"); + if (opcode == 0x887 && index > 0 && condition == 0 && insn.words[1] <= 0x7ffff && + code[index - 1].Bits(0, 12) == 0x88c && + code[index - 1].Bits(24, 8) == insn.Bits(24, 8) && + code[index - 1].words[1] == insn.words[1] && code[index - 1].words[2] == 0x0bf01070 && + insn.words[2] == 0x08000000) { + // Preserve the unsigned clamp idiom for indirect-table analysis. + emit(0x3820000000070000ULL | target | (u64{a} << 8) | + (u64{insn.words[1]} << 20) | (u64{7} << 39)); + break; + } + if (opcode == 0x887) mov_imm(252, insn.words[1]); + const u32 temporary = 6 - condition, mask = 1U << temporary; + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); + emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{mask} << 20)); + emit(0x5ca0000000070000ULL | target | (u64{a} << 8) | (u64{b} << 20) | + (u64{temporary} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); + break; + } + case 0x386: case 0x986: { // Global store of 32/64/128 bits, 64-bit address. + const 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) + return fail(index, "native global store size"); + u32 address = insn.Bits(24, 8); + bool extended = insn.Bits(72, 1) != 0; + if (opcode == 0x986) { + const u32 uniform_address = uniform_register(insn.Bits(64, 6)); + if ((insn.words[2] & ~(0x04000000U | 0x3fU | 0xe00U)) != 0x0810a100 || + uniform_address == ~0U || (uniform_address & 1) || !extended) + return fail(index, "native global store uniform addressing mode"); + if (address == 255) address = uniform_address; + else { + if (insn.Bits(90, 1) && ((address & 1) || address >= 254)) + return fail(index, "native global store 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; + } + } else if ((insn.words[2] & ~0xe00U) != 0x0010a100) { + return fail(index, "native global store register addressing mode"); + } + emit(0xeed8000000000000ULL | predicate | data | + (u64{address} << 8) | (u64{insn.Bits(40, 24)} << 20) | + (u64{extended} << 45) | (u64{size} << 48)); + break; + } + case 0x95b: // Fragment discard; execution predicate is preserved. + if (insn.words[1] || insn.words[2] != 0x03800000) + return fail(index, "native fragment discard mode"); + emit(0xe33000000000000fULL | predicate); + break; + case 0x305: case 0xb05: { // F2I NTZ, register/constant-buffer source + const bool half = opcode == 0x305 && insn.words[2] == 0x0010f000; + const bool small = insn.words[2] == 0x00212800; + const bool floor = (insn.words[2] & ~0x100U) == 0x00207000; + const bool nearest = (insn.words[2] & ~0x100U) == 0x00203000; + if ((!half && !small && !floor && !nearest && (insn.words[2] & ~0x100U) != 0x0021f000) || (opcode == 0x305 && (insn.words[1] & 0xffffff00U))) { + return fail(index, "native float-to-integer conversion mode"); + } + const bool is_signed = insn.Bits(72, 1); + u32 source = opcode == 0xb05 ? materialize(2, 0, 252, false, false) : insn.Bits(32, 8); + if (half) { + extract_half({source, 2, false, false}, 0, 252); + source = 252; + } + if (!is_signed) { + // Clear NaNs without consuming a guest predicate. Clamp the + // conversion below 2^32, retaining an overflow mask for saturation. + emit(0x5800000000000000ULL | predicate | 250 | (u64{source} << 8) | + (u64{source} << 20) | (u64{7} << 39) | (u64{7} << 48)); + emit(0x5c40000000000000ULL | predicate | 250 | (u64{source} << 8) | + (u64{250} << 20) | (u64{7} << 48)); + if (!small) { + mov_imm(252, 0x4f800000); + emit(0x5800000000000000ULL | predicate | 251 | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39) | (u64{6} << 48)); + mov_imm(252, 0x4f7fffff); + emit(0x5c60000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39)); + } + source = 250; + } + emit(base(0x5cb0000000000000ULL) | (u64{small ? 1U : 2U} << 8) | (u64{2} << 10) | + (u64{is_signed} << 12) | (u64{source} << 20) | + (u64{small || nearest ? 0U : floor ? 1U : 3U} << 39) | (u64{insn.Bits(80, 1)} << 44)); + if (!is_signed && !small) { + emit(base(0x5c40000000000000ULL) | (dest << 8) | (u64{251} << 20) | + (u64{1} << 41) | (u64{7} << 48)); + } + break; + } + case 0xb6f: { // TXQ texture dimensions, two packed scalar outputs. + if (insn.Bits(59, 5) || insn.Bits(64, 8) != 255 || insn.Bits(72, 4) != 3 || + insn.Bits(76, 14) || insn.Bits(91, 1) || insn.Bits(54, 5) != 2 || + insn.Bits(40, 14) > 0x1fff || (dest & 1) || dest >= 254) + return fail(index, "native texture dimension query mode"); + result.required_texture_buffer = 2; + emit(0xdf48000000000000ULL | predicate | dest | (u64{insn.Bits(24, 8)} << 8) | + (u64{1} << 22) | (u64{3} << 31) | (u64{insn.Bits(40, 14)} << 36) | + (u64{insn.Bits(90, 1)} << 49)); + break; + } + case 0xb66: { // TLD.SCR.LZ, direct texel fetch + const u32 mask = insn.Bits(72, 4); + const u32 dimension = insn.Bits(61, 3); + const bool array = dimension == 5; + const bool volume = dimension == 2; + u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); + const u32 second_dest = insn.Bits(64, 8); + const bool split = second_dest != 255; + constexpr std::array<u32, 8> rg_masks{1, 2, 4, 8, 3, 9, 10, 12}; + constexpr std::array<u32, 5> rgba_masks{7, 11, 13, 14, 15}; + u32 swizzle = ~0U; + if (split) { + for (u32 i = 0; i < rgba_masks.size(); ++i) if (rgba_masks[i] == mask) swizzle = i; + } else { + for (u32 i = 0; i < rg_masks.size(); ++i) if (rg_masks[i] == mask) swizzle = i; + } + if (!insn.Bits(60, 1) || (dimension != 0 && dimension != 1 && !array && !volume) || + swizzle == ~0U || insn.Bits(76, 5) || + insn.Bits(81, 3) != 7 || insn.Bits(87, 3) != 1 || + insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) >= 18 || + dest == 255 || (mask & (mask - 1) && ((dest & 1) || dest >= 254)) || + (split && mask == 15 && ((second_dest & 1) || second_dest >= 254)) || + (array && ((source_a & 1) || source_a >= 249)) || + (volume && ((source_a & 1) || source_a >= 254))) { + return fail(index, "native texel fetch mode or descriptor binding"); + } + const u32 binding = insn.Bits(54, 5); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return fail(index, "native texture descriptor buffer conflict"); + result.required_texture_buffer = binding; + if (array) { + emit(0x5c98000000000000ULL | predicate | 250 | (u64{15} << 39) | (u64{source_a + 1} << 20)); + emit(0x5c98000000000000ULL | predicate | 251 | (u64{15} << 39) | (u64{source_b} << 20)); + source_b = 250; + } + emit(0xda00000000000000ULL | predicate | dest | (u64{source_a} << 8) | + (u64{source_b} << 20) | (u64{insn.Bits(64, 8)} << 28) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{swizzle} << 50) | (u64{array ? 8U : volume ? 7U : dimension == 0 ? 0U : 2U} << 53)); + break; + } + case 0x91a: // DEPBAR.LE SB5, 1: scoreboard wait, represented by IR data dependencies. + if (insn.words[0] != 0x791a || (insn.words[1] & ~0xfc0U) != 0xd000 || insn.words[2]) { + return fail(index, "native dependency barrier mode"); + } + break; + case 0xb60: { // TEX, scalar 2D implicit-LOD sample + const u32 mask = insn.Bits(72, 4); + const u32 second_dest = insn.Bits(64, 8); + constexpr std::array<u32, 8> rg_masks{1, 2, 4, 8, 3, 9, 10, 12}; + constexpr std::array<u32, 5> rgba_masks{7, 11, 13, 14, 15}; + u32 swizzle = ~0U; + const bool split = second_dest != 255; + const bool depth_compare = insn.Bits(78, 1); + const u32 dimension = insn.Bits(61, 3), lod = insn.Bits(87, 3); + const u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); + if (dimension == 5 && depth_compare && lod == 0 && !split && mask == 1) { + if (!insn.Bits(60, 1) || insn.Bits(76, 2) || insn.Bits(81, 3) != 7 || + insn.Bits(59, 1) || insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) != 2 || + (source_a & 1) || (source_b & 1) || source_a >= 248 || source_b >= 248 || dest >= 248) + return fail(index, "native array depth comparison sample mode"); + result.required_texture_buffer = 2; + for (u32 i = 0; i < 3; ++i) { + const u32 source = i < 2 ? source_a + i : source_b; + emit(0x5c98000000000000ULL | predicate | (250 + i) | + (u64{15} << 39) | (u64{source} << 20)); + } + emit(0xc000000000000000ULL | predicate | dest | (u64{250} << 8) | + (u64{source_b + 1} << 20) | (u64{3} << 28) | (u64{1} << 31) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{1} << 50) | (u64{7} << 51)); + break; + } + // General form: gather the native operand list (array index first, coordinates, + // then LOD/bias and the depth reference) into Maxwell TEX's coordinate and meta + // registers, sample into scratch and copy the packed components out. Native packs + // up to four values as halves of Ra and Rb (1+1, 2+1, 2+2), five or more as 4+n. + const auto general_tex = [&]() -> bool { + constexpr std::array<u32, 8> coordinate_count{1, 2, 3, 3, 2, 3, 0, 4}; + constexpr std::array<u32, 8> 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) || + 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 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; + return source_b == 255 ? 255U : source_b + (i - in_a); + }; + if ((source_a != 255 && source_a + in_a > 246) || + (count > in_a && source_b != 255 && source_b + (count - in_a) > 246)) + return false; + const u32 binding = insn.Bits(54, 5); + if (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding) + return false; + result.required_texture_buffer = binding; + 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. + 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))); + 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)); + const u32 components = static_cast<u32>(std::popcount(mask)); + for (u32 i = 0; i < components; ++i) { + u32 target = 255; + if (!split) target = dest == 255 ? 255U : static_cast<u32>(dest) + i; + else if (i < 2) target = dest == 255 ? 255U : static_cast<u32>(dest) + i; + else target = second_dest + (i - 2); + if (target != 255) copy(target, 250 + i); + } + return true; + }; + if (split) { + for (u32 i = 0; i < rgba_masks.size(); ++i) if (rgba_masks[i] == mask) swizzle = i; + } else { + for (u32 i = 0; i < rg_masks.size(); ++i) if (rg_masks[i] == mask) swizzle = i; + } + // SM70 scalar coordinates and split RG/BA destinations match TEXS. + // Descriptor offsets are words in both encodings, but only native + // explicitly encodes the constant-buffer binding. + if (!insn.Bits(60, 1) || (dimension != 1 && dimension != 5) || + swizzle == ~0U || insn.Bits(76, 2) || + insn.Bits(81, 3) != 7 || (lod > 1 && (lod != 3 || dimension != 1 || depth_compare)) || + (lod == 3 && ((source_a & 1) || source_a >= 249)) || + insn.Bits(59, 1) || + insn.Bits(40, 14) > 0x1fff || insn.Bits(54, 5) != 2 || + dest == 255 || (mask & (mask - 1) && ((dest & 1) || dest >= 254)) || + (split && mask == 15 && ((second_dest & 1) || second_dest >= 254)) || + (depth_compare && (dimension != 1 || split || mask != 1 || (source_a & 1) || source_a >= 249)) || + (dimension == 5 && ((source_a & 1) || source_a >= 249))) { + if (!general_tex()) return fail(index, "native texture mode or descriptor binding"); + break; + } + result.required_texture_buffer = 2; + emit(0xd800000000000000ULL | predicate | dest | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(64, 8)} << 28) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | + (u64{swizzle} << 50) | (u64{dimension == 5 ? (lod ? 8U : 7U) : depth_compare ? (lod ? 6U : 4U) : (lod == 3 ? 3U : lod ? 2U : 1U)} << 53)); + break; + } + case 0x816: { // Captured PRMT half-word concatenation + if ((insn.words[1] != 0x5410 && insn.words[1] != 0x5432 && insn.words[1] != 0x7610 && insn.words[1] != 0x7632) || + (insn.words[2] & ~0xffU)) return fail(index, "native permutation mode"); + const u32 a = insn.Bits(24, 8), b = insn.Bits(64, 8); + if (insn.words[1] != 0x5432 && insn.words[1] != 0x7632) { + emit(0x0400000000000000ULL | predicate | 250 | (u64{a} << 8) | (u64{0xffff} << 20)); + } else { + emit(0x3828000000000000ULL | predicate | 250 | (u64{a} << 8) | (u64{16} << 20)); + } + if (insn.words[1] == 0x7610 || insn.words[1] == 0x7632) { + emit(0x0400000000000000ULL | predicate | 251 | (u64{b} << 8) | (u64{0xffff0000} << 20)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{251} << 20)); + } else { + emit(base(0x5c18000000000000ULL) | (u64{b} << 8) | (u64{250} << 20) | (u64{16} << 39)); + } + break; + } + case 0x307: { // FRND.F32, round to an integral floating-point value + if (insn.Bits(75, 2) != 2 || insn.Bits(84, 2) != 2 || insn.Bits(81, 3) || + insn.Bits(86, 6)) return fail(index, "native floating round format"); + emit(base(0x5ca8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{8 + insn.Bits(78, 2)} << 39) | + (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(63, 1)} << 45) | + (u64{insn.Bits(62, 1)} << 49)); + break; + } + case 0x321: // ALD + case 0x322: { // AST + if (insn.Bits(76, 2) || insn.Bits(79, 1)) return fail(index, "native indexed/patch attribute access"); + const u32 address = insn.Bits(40, 10); + const u32 components = insn.Bits(74, 2) + 1; + 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<u32>(dest) : insn.Bits(32, 8); + 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) | + predicate | value_reg | (u64{255} << 8) | (u64{address} << 20) | + (u64{vertex} << 39) | (u64{components - 1} << 47); + result.instructions.push_back(translated); + break; + } + case 0x387: case 0x983: { // Captured scalar local spill at byte offset zero. + const bool store = opcode == 0x387; + if (insn.Bits(24, 8) != 255 || (store && (insn.Bits(16, 8) || insn.words[1] > 247)) || + (!store && (insn.words[1] || dest >= 248)) || + insn.words[2] != (store ? 0x00100800U : 0x00300800U)) + return fail(index, "native local scalar spill mode"); + emit((store ? 0xef50000000000000ULL : 0xef40000000000000ULL) | predicate | + (store ? insn.words[1] : dest) | (u64{255} << 8) | (u64{4} << 48)); + break; + } + case 0x803: { // P2R PR, register merge source, immediate mask, byte select (76-77) + if (insn.words[1] > 0x7f || (insn.words[2] & ~0x3000U)) + return fail(index, "native predicate-to-register mode"); + emit(base(0x38e8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.words[1]} << 20) | (u64{insn.Bits(76, 2)} << 41)); + break; + } + case 0x804: { // R2P PR, register, immediate mask, byte select (76-77) + if (insn.Bits(16, 8) || insn.words[1] > 0x7f || (insn.words[2] & ~0x3000U)) + return fail(index, "native register-to-predicate mode"); + emit(0x38f0000000000000ULL | predicate | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.words[1]} << 20) | (u64{insn.Bits(76, 2)} << 41)); + break; + } + case 0xc02: // MOV uniform register + case 0x202: // MOV register + case 0x802: // MOV immediate + case 0xa02: { // MOV constant buffer + if (insn.Bits(72, 4) != 15 || (insn.Bits(76, 16) & ~(opcode == 0xc02 ? 0x8000U : 0U))) return fail(index, "native MOV modifiers"); + if (opcode == 0x802) { + result.instructions.push_back(base(0x0100000000000000ULL) | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); + } else if (opcode == 0xc02) { + if (insn.words[1] >= 16 || insn.words[2] != 0x08000f00) return fail(index, "native uniform move source"); + emit(base(0x5c98000000000000ULL) | (u64{15} << 39) | (u64{insn.words[1] + 216} << 20)); + } else if (opcode == 0x202) { + if (insn.Bits(40, 24)) return fail(index, "native MOV register modifiers"); + result.instructions.push_back(base(0x5c98000000000000ULL) | (u64{15} << 39) | + (u64{insn.Bits(32, 8)} << 20)); + } else { + if (!valid_cbuf() || insn.Bits(62, 2)) return fail(index, "native MOV constant buffer"); + result.instructions.push_back(base(0x4c98000000000000ULL) | (u64{15} << 39) | cbuf()); + } + break; + } + case 0x223: case 0x423: case 0x623: case 0x823: case 0xa23: case 0xc23: case 0xe23: { // FFMA + if (insn.Bits(76, 1) || (insn.Bits(81, 11) & ~(opcode == 0xe23 || opcode == 0xc23 ? 0x400U : 0U))) return fail(index, "native FFMA denormal/source mode"); + const u32 a = operand(0), b = operand(1), c = operand(2); + emit(base(0x5980000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | (u64{c} << 39) | + (u64{insn.Bits(77, 1)} << 50) | (u64{insn.Bits(78, 2)} << 51) | (u64{insn.Bits(80, 1)} << 53)); + break; + } + case 0x220: case 0x820: case 0xa20: case 0xc20: { // FMUL + const u32 scale = insn.Bits(84, 3); + if (insn.Bits(76, 1) || scale == 0 || insn.Bits(81, 3) || + insn.Bits(87, 4) || insn.Bits(91, 1) != (opcode == 0xc20) || + (scale != 4 && (!insn.Bits(80, 1) || insn.Bits(78, 2)))) + return fail(index, "native FMUL scale/mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5c68000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(78, 2)} << 39) | (u64{std::array<u32, 8>{0, 3, 2, 1, 0, 6, 5, 4}[scale]} << 41) | (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(77, 1)} << 50)); + break; + } + case 0x221: case 0x421: case 0x621: case 0x821: case 0xa21: case 0xe21: { // FADD + if (insn.Bits(81, 10) || insn.Bits(91, 1) != (opcode == 0xe21)) return fail(index, "native FADD mode"); + const u32 a = operand(0), b = operand(insn.Bits(9, 3) == 2 || insn.Bits(9, 3) == 3 || insn.Bits(9, 3) == 7 ? 2 : 1); + emit(base(0x5c58000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(78, 2)} << 39) | (u64{insn.Bits(80, 1)} << 44) | (u64{insn.Bits(77, 1)} << 50)); + break; + } + case 0x209: case 0x809: case 0xa09: { // FMNMX + if (insn.Bits(74, 6) || insn.Bits(81, 6) || insn.Bits(91, 1)) + return fail(index, "native floating min/max mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5c60000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42) | + (u64{insn.Bits(80, 1)} << 44)); + break; + } + case 0x308: case 0xb08: { // MUFU + const u32 operation = insn.Bits(74, 6); + if (operation > 8 || operation == 6 || operation == 7 || insn.Bits(80, 12)) return fail(index, "native MUFU mode"); + if (insn.Bits(73, 1)) { + if (opcode != 0x308 || insn.Bits(60, 2) > 1) + return fail(index, "native half special-function source"); + emit(0x5ca8000000000000ULL | predicate | 250 | (u64{2} << 8) | (u64{1} << 10) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(60, 1)} << 41) | + (u64{insn.Bits(63, 1)} << 45) | (u64{insn.Bits(62, 1)} << 49)); + emit(0x5080000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{operation} << 20)); + emit(base(0x5ca8000000000000ULL) | (u64{1} << 8) | (u64{2} << 10) | (u64{250} << 20)); + } else { + const u32 src = operand(1); + emit(base(0x5080000000000000ULL) | (u64{src} << 8) | (u64{operation} << 20)); + } + break; + } + case 0x526: { // IPA.PASS / CONSTANT with the attribute indexed by a uniform register + const u32 index_register = uniform_register(insn.Bits(24, 8)); + if (index_register == ~0U || insn.Bits(32, 8) != 255 || + (insn.Bits(76, 4) != 0 && insn.Bits(76, 4) != 4) || insn.Bits(81, 3) != 7) + return fail(index, "native indexed IPA mode"); + if (index_register == 255) { + emit(base(0xe000000000000000ULL) | (u64{255} << 8) | (u64{255} << 20) | + (u64{insn.Bits(64, 8)} << 30) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + // Maxwell IPA.IDX reads the whole byte address from the index register. + emit(0x1c00000000000000ULL | predicate | 250 | (u64{index_register} << 8) | + (u64{insn.Bits(64, 8) * 4} << 20)); + emit(base(0xe000000000000000ULL) | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(64, 8)} << 30) | (u64{1} << 38) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + case 0x326: { // IPA.PASS / CONSTANT, default sample location + if (insn.Bits(32, 8) != 255 || (insn.Bits(76, 4) != 0 && insn.Bits(76, 4) != 4) || insn.Bits(81, 3) != 7) return fail(index, "native IPA interpolation mode"); + emit(base(0xe000000000000000ULL) | (u64{255} << 8) | (u64{255} << 20) | (u64{insn.Bits(64, 8)} << 30) | (u64{insn.Bits(78, 1)} << 55)); + break; + } + case 0x805: { // CS2R SRZ + if (insn.Bits(72, 8) != 255 || !insn.Bits(80, 1) || dest >= 249) return fail(index, "native special-register read"); + mov_imm(static_cast<u32>(dest), 0); + mov_imm(static_cast<u32>(dest) + 1, 0); + break; + } + case 0x306: { // I2F.F32.U8, low byte, round-to-nearest + if (insn.words[2] != 0x1000 || (insn.words[1] & 0xffffff00U)) + return fail(index, "native byte-to-float conversion mode"); + emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{insn.Bits(32, 8)} << 20)); + 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"); + const u32 src = operand(1, false); + emit(base(0x5cb8000000000000ULL) | (u64{2} << 8) | (u64{2} << 10) | (u64{insn.Bits(74, 1)} << 13) | (u64{src} << 20)); + break; + } + case 0x207: case 0x807: case 0xa07: + case 0x208: case 0x808: case 0xa08: { // Bitwise SEL / FSEL + const u32 a = operand(0, false), b = operand(1, false); + emit(base(0x5ca0000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + break; + } + case 0x20a: case 0x80a: case 0xa0a: { // FSET.BF with PT/AND + if (insn.Bits(64, 8) || insn.Bits(74, 2) || insn.Bits(81, 6) || + insn.Bits(87, 4) != 7 || insn.Bits(91, 1)) + return fail(index, "native floating comparison output mode"); + const u32 a = operand(0), b = operand(1); + emit(base(0x5800000000000000ULL) | (u64{a} << 8) | (u64{b} << 20) | + (u64{7} << 39) | (u64{insn.Bits(76, 4)} << 48) | (u64{1} << 52) | + (u64{insn.Bits(80, 1)} << 55)); + break; + } + case 0x20b: case 0x80b: case 0xa0b: // FSETP + case 0x20c: case 0x80c: case 0xa0c: case 0xc0c: { // ISETP + const bool floating = (opcode & 0x1ff) == 0xb; + if (!floating && insn.Bits(72, 1)) return fail(index, "native extended integer comparison"); + const u32 a = operand(0, floating), b = operand(1, floating); + const u32 combine = insn.Bits(74, 2); + if (combine > 2) return fail(index, "native predicate combination"); + u64 out = (floating ? 0x5bb0000000000000ULL : 0x5b60000000000000ULL) | predicate | + insn.Bits(84, 3) | (u64{insn.Bits(81, 3)} << 3) | (u64{a} << 8) | (u64{b} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42) | (u64{combine} << 45); + if (floating) out |= (u64{insn.Bits(80, 1)} << 47) | (u64{insn.Bits(76, 4)} << 48); + else out |= (u64{insn.Bits(73, 1)} << 48) | (u64{insn.Bits(76, 3)} << 49); + emit(out); + break; + } + case 0x212: case 0x812: case 0xc12: { // Bounded LOP3 bitwise forms, no predicate output + 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) { + 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<u32>(dest), 250, 252, 1, pred_target); + } else if (lut == 0xf8) { + logical(250, b, c, false); + logical(static_cast<u32>(dest), a, 250, 1, pred_target); + } else if (lut == 0x48) { + logical(250, a, c, 2); + logical(static_cast<u32>(dest), 250, b, 0, pred_target); + } else { + logical(static_cast<u32>(dest), a, b, lut == 0xfc ? 1U : 0U, pred_target); + } + break; + } + case 0x217: case 0x817: case 0xa17: case 0xc17: { // IMNMX register, immediate, cbuf or uniform operand. + if ((insn.words[2] & ~(0x04000200U | (opcode == 0xc17 ? 0x08000000U : 0U))) != 0x03800000 || + (opcode != 0xa17 && insn.words[1] > (opcode == 0x817 ? 0x7ffffU : 255U)) || + (opcode == 0xa17 && (!valid_cbuf() || insn.Bits(62, 2)))) + return fail(index, "native integer min/max mode"); + const u32 source = opcode == 0xc17 ? uniform_register(insn.words[1]) + : opcode == 0xa17 ? operand(1, false) + : insn.words[1]; + if (source == ~0U) return fail(index, "native uniform min/max operand"); + emit(base(opcode == 0x817 ? 0x3820000000000000ULL : 0x5c20000000000000ULL) | + (u64{insn.Bits(24, 8)} << 8) | (u64{source} << 20) | + (u64{insn.Bits(87, 4)} << 39) | (u64{insn.Bits(73, 1)} << 48)); + break; + } + case 0x819: { // SHF.HI, immediate shift + const bool is_signed = insn.Bits(73, 2) == 2; + if (!insn.Bits(80, 1) && !insn.Bits(76, 1) && insn.Bits(64, 8) == 255 && + insn.Bits(73, 2) == 3 && !insn.Bits(75, 1) && insn.words[1] <= 32) { + emit(base(0x3848000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{insn.words[1]} << 20)); + break; + } + if ((insn.Bits(73, 2) != 3 && !is_signed) || insn.Bits(75, 1) || !insn.Bits(80, 1) || insn.words[1] > 32) return fail(index, "native funnel-shift mode"); + const u32 low = insn.Bits(24, 8), high = insn.Bits(64, 8), shift = insn.words[1]; + if (insn.Bits(76, 1)) { + if (low != 255) return fail(index, "native nonzero low right-shift word"); + emit(base(0x3828000000000000ULL) | (u64{high} << 8) | (u64{shift} << 20) | + (u64{is_signed} << 48)); + } else { + if (is_signed) return fail(index, "native signed left funnel shift"); + emit(base(0x36f8000000000000ULL) | (u64{low} << 8) | (u64{high} << 39) | (u64{shift} << 20)); + } + break; + } + case 0xa11: { // LEA with constant base and carry predicate. + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(80, 1) || + insn.Bits(81, 3) == 7 || insn.Bits(62, 2)) + return fail(index, "native constant LEA carry mode"); + const u32 b = operand(1, false); + emit(base(0x5c18000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(75, 5)} << 39)); + emit(0x5b60000000000000ULL | predicate | 7 | (u64{insn.Bits(81, 3)} << 3) | + (dest << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{1} << 49)); + break; + } + case 0x211: case 0x811: { // LEA low word, no carry result + if (opcode == 0x211 && insn.Bits(80, 1)) { + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(81, 3) != 7 || + insn.Bits(32, 8) != 255 || insn.Bits(64, 8) != 255 || insn.Bits(62, 2)) + return fail(index, "native high effective address mode"); + emit(base(0x3828000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{32 - insn.Bits(75, 5)} << 20)); + break; + } + if (insn.Bits(72, 1) || insn.Bits(74, 1) || insn.Bits(80, 1) || insn.Bits(81, 3) != 7 || + (opcode == 0x211 && insn.Bits(62, 2))) return fail(index, "native LEA carry/high mode"); + u32 b = insn.Bits(32, 8); + if (opcode == 0x811) { mov_imm(251, insn.words[1]); b = 251; } + emit(base(0x5c18000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(75, 5)} << 39)); + break; + } + case 0x210: case 0x810: case 0xa10: case 0xc10: { // IADD3 with zero third source and no carry outputs + if (opcode == 0xa10 && insn.Bits(81, 3) != 7) { + if (insn.Bits(64, 8) != 255 || insn.Bits(73, 4) || insn.Bits(84, 7) != 0x7f || + insn.Bits(77, 4) != 15 || insn.Bits(63, 1) || insn.Bits(72, 1)) + return fail(index, "native constant add carry mode"); + const u32 b = operand(1, false); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{b} << 20)); + emit(0x5b60000000000000ULL | predicate | 7 | (u64{insn.Bits(81, 3)} << 3) | + (dest << 8) | (u64{b} << 20) | (u64{7} << 39) | (u64{1} << 49)); + break; + } + if (opcode == 0xa10 && insn.Bits(74, 1)) { + if (insn.Bits(24, 8) != 255 || (insn.words[2] & ~0x07800000U) != 0x007fe4ffU) + return fail(index, "native extended constant add mode"); + const u32 b = operand(1, false); + mov_imm(250, 1); + emit(0x5ca0000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{255} << 20) | + (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); + emit(base(0x5c10000000000000ULL) | (u64{b} << 8) | (u64{250} << 20)); + break; + } + if (insn.Bits(64, 8) != 255 && !insn.Bits(73, 4) && insn.Bits(77, 7) == 0x7f && insn.Bits(84, 7) == 0x7f) { + const u32 b = operand(1, false); + emit(0x5c10000000000000ULL | predicate | 250 | (u64{insn.Bits(24, 8)} << 8) | + (u64{b} << 20) | (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{insn.Bits(64, 8)} << 20)); + break; + } + if (insn.Bits(64, 8) != 255 || insn.Bits(73, 4) || insn.Bits(77, 7) != 0x7f || insn.Bits(84, 7) != 0x7f) return fail(index, "native integer carry/third source"); + if (opcode == 0xa10 || opcode == 0xc10) { + if (opcode == 0xc10 && (insn.words[1] & ~0x8000003fU)) return fail(index, "native uniform add operand"); + const u32 b = operand(1, false); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{b} << 20) | + (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + } else if (opcode == 0x210) { + if (insn.words[1] & ~0x800000ffU) return fail(index, "native integer add register modifiers"); + emit(base(0x5c10000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(63, 1)} << 48) | (u64{insn.Bits(72, 1)} << 49)); + } else { + emit(base(0x1c00000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | (u64{insn.words[1]} << 20) | (u64{insn.Bits(72, 1)} << 56)); + } + break; + } + case 0x949: case 0x958: { // Captured BRX / BRXU with absolute native-PC table values + if (insn.Bits(12, 4) != 7 || (insn.words[2] & ~0x3ffffU) != (opcode == 0x958 ? 0x0b800000U : 0x03800000U)) + return fail(index, "native computed branch mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast<s64>(raw << 14) >> 14; + if (s64{index} * 16 + 16 + relative != 0) + return fail(index, "native computed branch relative base"); + const s64 offset = -s64{address_of(static_cast<u32>(result.instructions.size()))} - 8; + if (offset < -0x800000) return fail(index, "native computed branch range"); + const u32 source = opcode == 0x958 ? uniform_register(insn.Bits(24, 8)) : insn.Bits(24, 8); + if (source == ~0U) return fail(index, "native uniform computed branch source"); + emit(0xe25000000000000fULL | predicate | (u64{source} << 8) | + ((static_cast<u64>(offset) & 0xffffff) << 20)); + break; + } + case 0x944: { // Verified CALL.REL.NOINC with R55 holding the return PC + // A predicated NOINC forward transfer can terminate a loop at its + // reconvergence point without a return-link setup or matching RET. + if (native_predicate != 7 && insn.words[1] == 16 && insn.words[2] == 0x03c00000 && + index + 2 < code.size() && code[index + 1].Bits(0, 12) == 0x947) { + relocations.push_back({static_cast<u32>(result.instructions.size()), index + 2}); + emit(0xe24000000000000fULL | predicate); + break; + } + if (main_exited || index == 0 || insn.words[0] != 0x7944 || + (insn.words[2] & ~0x3ffffU) != 0x03c00000 || + (code[index - 1].words[0] & ~0x00ff0000U) != 0x00007802 || + code[index - 1].Bits(16, 8) >= 250 || + code[index - 1].words[1] != (index + 1) * 16 || + code[index - 1].words[2] != 0xf00) + return fail(index, "native call linkage mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast<s64>(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16) return fail(index, "native call destination alignment"); + const u32 link = code[index - 1].Bits(16, 8); + for (size_t call = 0; call < call_targets.size(); ++call) { + if (call_targets[call] == target / 16 && call_links[call] != link) + return fail(index, "native inconsistent call linkage"); + } + call_targets.push_back(static_cast<u32>(target / 16)); + call_links.push_back(link); + relocations.push_back({static_cast<u32>(result.instructions.size()), target / 16}); + emit(0xe26000000007000fULL); + break; + } + case 0x950: { // Verified RET.REL.NODEC epilogue restoring the call link + const u32 return_reg = insn.Bits(24, 8); + if (!main_exited || index < 2 || !active_link || + (insn.words[0] & 0x00ffffffU) != 0x00007950 || return_reg >= 249 || + (insn.words[2] & ~0x3ffffU) != 0x03c00000 || + code[index - 2].words[0] != (0x00007202U | (return_reg << 16)) || + code[index - 2].words[1] != *active_link || code[index - 2].words[2] != 0xf00 || + code[index - 1].words[0] != (0x00007802U | ((return_reg + 1) << 16)) || + code[index - 1].words[1] || code[index - 1].words[2] != 0xf00 || !barrier_tokens.empty()) + return fail(index, "native return linkage mode"); + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast<s64>(raw << 14) >> 14; + if (s64{index} * 16 + 16 + relative != 0) + return fail(index, "native return relative base"); + for (const auto& relocation : relocations) { + if (relocation.target == index || relocation.target == index - 1) + return fail(index, "native branch bypasses return linkage"); + } + emit(0xe32000000007000fULL); + bool complete = true; + for (u32 target : call_targets) complete &= target <= index; + if (complete) return finish(index); + break; + } + case 0x945: case 0x941: case 0x942: case 0x948: // (+ WARPSYNC) + // BSSY / BSYNC / BREAK only manage warp reconvergence: BSYNC waits for the other + // threads and falls through, BREAK drops the thread from that wait. Eden runs threads + // independently, so all three are no-ops (a SYNC would wrongly jump, and a branch + // out of a lowered SSY region would leave its token on Eden's flow stack). + break; + case 0x947: { // BRA + u64 branch_predicate = predicate; + if (opcode == 0x947 && insn.Bits(87, 4) != 7) { + if (insn.Bits(12, 4) != 7) { + relocations.push_back({static_cast<u32>(result.instructions.size()), index + 1}); + emit(0xe24000000000000fULL | (u64{insn.Bits(12, 4) ^ 8U} << 16)); + } + branch_predicate = u64{insn.Bits(87, 4)} << 16; + } else if (insn.Bits(87, 4) != 7) { + return fail(index, "native extra branch condition"); + } + const u64 raw = insn.words[1] | (u64{insn.words[2] & 0x3ffff} << 32); + const s64 relative = static_cast<s64>(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16) return fail(index, "native branch destination alignment"); + relocations.push_back({static_cast<u32>(result.instructions.size()), target / 16}); + emit(0xe24000000000000fULL | branch_predicate); + break; + } + case 0x94d: { // EXIT + if (insn.Bits(84, 2) || insn.Bits(87, 4) != 7) { + return fail(index, "native conditional/special EXIT"); + } + result.instructions.push_back(0xe30000000000000fULL | predicate); + if (insn.Bits(12, 4) != 7) break; + if (!barrier_tokens.empty()) return fail(index, "native unclosed reconvergence scope"); + 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); + break; + } + if (main_exited) return fail(index, "native library exits without returning"); + for (u32 target : call_targets) { + if (target <= index) return fail(index, "native call target overlaps main program"); + } + main_exited = true; + break; + } + default: + return fail(index, "unsupported native opcode"); + } + if (!operand_error.empty()) return fail(index, operand_error); + if (guarded_uniform) restore_predicate(); + } + return fail(static_cast<u32>(code.size()), main_exited ? "native program lacks captured library return" : "native program has no unconditional EXIT"); +} + +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_address_map.h b/src/shader_recompiler/frontend/ampere/program_address_map.h new file mode 100644 index 0000000000..bee32a821c --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_address_map.h @@ -0,0 +1,29 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once +#include <optional> +#include <span> +#include <vector> +#include "common/common_types.h" + +namespace Shader::Ampere { +constexpr u64 NativeAddressMapMagic = 0x50414d4e4544454eULL; +inline bool HasNativeAddressMap(std::span<const u64> code) { + return code.size() >= 2 && code.back() == NativeAddressMapMagic; +} +inline void AppendNativeAddressMap(std::vector<u64>& code, std::span<const u32> offsets) { + if (offsets.empty()) return; + for (u32 index : offsets) code.push_back(0x50U + index * 8 + ((index + 1) / 3) * 8); + code.push_back(offsets.size()); + code.push_back(NativeAddressMapMagic); +} +inline std::optional<u32> NativeProgramAddress(std::span<const u64> code, u32 address) { + if (!HasNativeAddressMap(code) || address % 16) return {}; + const u64 count = code[code.size() - 2]; + if (!count || count > 65536 || count + 2 > code.size() || address / 16 >= count) return {}; + const size_t start = code.size() - static_cast<size_t>(count) - 2; + const u64 target = code[start + address / 16]; + if (target < 0x50 || target >= start * 8 || target % 8 || target % 32 == 0) return {}; + return static_cast<u32>(target); +} +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_capture.h b/src/shader_recompiler/frontend/ampere/program_capture.h new file mode 100644 index 0000000000..ea7dfb7ef6 --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_capture.h @@ -0,0 +1,43 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" + +namespace Shader::Ampere { +class ProgramCaptureState { +public: + enum class Result { Continue, Complete, InvalidTarget }; + static constexpr u32 MaximumInstructions = 65536; + + Result Observe(const Instruction& instruction, u32 index) { + const u32 opcode = instruction.Bits(0, 12); + if (opcode == 0x944 || opcode == 0x947 || opcode == 0x945) { + const u64 raw = instruction.words[1] | (u64{instruction.words[2] & 0x3ffff} << 32); + const s64 relative = opcode == 0x945 ? static_cast<s32>(instruction.words[1]) + : static_cast<s64>(raw << 14) >> 14; + const s64 target = s64{index} * 16 + 16 + relative; + if (target < 0 || target % 16 || target >= MaximumInstructions * 16) + return Result::InvalidTarget; + if (opcode == 0x944) { + last_call_target = std::max(last_call_target, static_cast<u32>(target / 16)); + } else { + last_forward_target = std::max(last_forward_target, static_cast<u32>(target / 16)); + } + } + if (opcode == 0x94d && instruction.Bits(12, 4) == 7 && index >= last_forward_target) { + if (!last_call_target) return Result::Complete; + main_exited = true; + } + if (opcode == 0x950 && instruction.Bits(12, 4) == 7 && main_exited && + index >= last_call_target && index >= last_forward_target) + return Result::Complete; + return Result::Continue; + } + +private: + u32 last_call_target{}; + u32 last_forward_target{}; + bool main_exited{}; +}; +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/ampere/program_override.h b/src/shader_recompiler/frontend/ampere/program_override.h new file mode 100644 index 0000000000..0943028b0c --- /dev/null +++ b/src/shader_recompiler/frontend/ampere/program_override.h @@ -0,0 +1,85 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include <cstring> +#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h" + +namespace Shader::Ampere { + +inline u64 NativeProgramKey(std::span<const Instruction> source) { + u64 hash = 14695981039346656037ULL; + for (const auto& instruction : source) { + for (u32 word : instruction.words) { + for (u32 byte = 0; byte < 4; ++byte) hash = (hash ^ ((word >> (byte * 8)) & 255)) * 1099511628211ULL; + } + } + return hash; +} + +// Private iteration format: version, lengths, texture bank, exact original +// instructions, then the newly generated canonical instructions. Source bytes +// are compared directly; the filename hash is only an index. +inline std::vector<u8> EncodeProgramOverride(std::span<const Instruction> source, + const LoweredProgram& program) { + if (!program || program.source_size != source.size_bytes() || program.source_instruction_offsets.size() != source.size()) return {}; + const std::array<u32, 6> header{0x4e454445, 0x3230564f, + static_cast<u32>(source.size_bytes()), static_cast<u32>(program.instructions.size()), + program.required_texture_buffer, static_cast<u32>(program.source_instruction_offsets.size())}; + std::vector<u8> bytes(sizeof(header) + source.size_bytes() + program.instructions.size() * 8 + program.source_instruction_offsets.size() * 4); + std::memcpy(bytes.data(), header.data(), sizeof(header)); + std::memcpy(bytes.data() + sizeof(header), source.data(), source.size_bytes()); + std::memcpy(bytes.data() + sizeof(header) + source.size_bytes(), program.instructions.data(), program.instructions.size() * 8); + std::memcpy(bytes.data() + sizeof(header) + source.size_bytes() + program.instructions.size() * 8, + program.source_instruction_offsets.data(), program.source_instruction_offsets.size() * 4); + return bytes; +} + +inline u32 ProgramOverrideSourceSize(std::span<const u8> bytes) { + std::array<u32, 6> header{}; + if (bytes.size() < sizeof(header)) return 0; + std::memcpy(header.data(), bytes.data(), sizeof(header)); + const bool mapped = header[1] == 0x3230564f; + if (header[0] != 0x4e454445 || (!mapped && header[1] != 0x3130564f) || + (mapped ? header[5] != header[2] / 16 : header[5] != 0) || + header[2] == 0 || header[2] % sizeof(Instruction) || header[2] > 1048576 || + header[3] == 0 || header[3] > 1048576 || (header[4] != ~0U && header[4] >= 18)) return 0; + if (bytes.size() != sizeof(header) + size_t{header[2]} + size_t{header[3]} * 8 + size_t{header[5]} * 4) return 0; + return header[2]; +} + +inline LoweredProgram DecodeProgramOverride(std::span<const u8> bytes, + std::span<const Instruction> source) { + LoweredProgram result; + const auto fail = [&result](std::string_view message) { + result.instructions.clear(); + result.source_instruction_offsets.clear(); + result.error = message; + return result; + }; + std::array<u32, 6> header{}; + if (!ProgramOverrideSourceSize(bytes)) return fail("invalid native override header or length"); + std::memcpy(header.data(), bytes.data(), sizeof(header)); + if (header[2] != source.size_bytes()) return fail("native override source length mismatch"); + if (std::memcmp(bytes.data() + sizeof(header), source.data(), source.size_bytes())) + return fail("native override source mismatch"); + result.instructions.resize(header[3]); + std::memcpy(result.instructions.data(), bytes.data() + sizeof(header) + source.size_bytes(), header[3] * 8); + if (header[5]) { + result.source_instruction_offsets.resize(header[5]); + std::memcpy(result.source_instruction_offsets.data(), + bytes.data() + sizeof(header) + source.size_bytes() + size_t{header[3]} * 8, + size_t{header[5]} * 4); + u32 previous = 0; + for (u32 offset : result.source_instruction_offsets) { + if (offset >= result.instructions.size() || offset < previous) + return fail("native override address map mismatch"); + previous = offset; + } + } + result.source_size = header[2]; + result.required_texture_buffer = header[4]; + return result; +} + +} // namespace Shader::Ampere diff --git a/src/shader_recompiler/frontend/maxwell/control_flow.cpp b/src/shader_recompiler/frontend/maxwell/control_flow.cpp index 7dabfbf16e..9fb30df664 100644 --- a/src/shader_recompiler/frontend/maxwell/control_flow.cpp +++ b/src/shader_recompiler/frontend/maxwell/control_flow.cpp @@ -457,7 +457,7 @@ CFG::AnalysisState CFG::AnalyzeBRX(Block* block, Location pc, Instruction inst, block->indirect_branches.reserve(targets.size()); for (const u32 target : targets) { - Block* const branch{AddLabel(block, block->stack, target, function_id)}; + Block* const branch{AddLabel(block, block->stack, env.TranslateProgramAddress(target), function_id)}; block->indirect_branches.push_back({ .block = branch, .address = target, diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt index 312a6a09db..d17deacacd 100644 --- a/src/video_core/CMakeLists.txt +++ b/src/video_core/CMakeLists.txt @@ -90,6 +90,7 @@ add_library(video_core STATIC host1x/vic.h macro.cpp macro.h + macro_turing.h fence_manager.h gpu.cpp gpu.h @@ -99,6 +100,8 @@ add_library(video_core STATIC invalidation_accumulator.h memory_manager.cpp memory_manager.h + native_qmd.h + native_query_trace.h nx2.h present.h pte_kind.h diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index f30c85911f..126e73c9b6 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -9,6 +9,7 @@ #include <algorithm> #include <memory> #include <numeric> +#include <stdexcept> #include "common/range_sets.inc" #include "video_core/buffer_cache/buffer_cache_base.h" @@ -1372,7 +1373,12 @@ void BufferCache<P>::UpdateVertexBuffer(u32 index) { const auto& array = maxwell3d->regs.vertex_streams[index]; const auto& limit = maxwell3d->regs.vertex_stream_limits[index]; const GPUVAddr gpu_addr_begin = array.Address(); - const GPUVAddr gpu_addr_end = limit.Address() + 1; + const bool native = maxwell3d->GetNativeMacroEngine().native_turing; + const u64 native_size = native ? maxwell3d->regs.NativeVertexStreamSize(index) : 0; + if (native && array.enable != 0 && native_size > (std::numeric_limits<u32>::max)()) { + throw std::runtime_error("Native vertex stream exceeds supported 32-bit buffer size"); + } + const GPUVAddr gpu_addr_end = native ? gpu_addr_begin + native_size : limit.Address() + 1; const std::optional<DAddr> device_addr = gpu_memory->GpuToCpuAddress(gpu_addr_begin); const u32 address_size = static_cast<u32>(gpu_addr_end - gpu_addr_begin); u32 size = address_size; // TODO: Analyze stride and number of vertices diff --git a/src/video_core/control/channel_state.cpp b/src/video_core/control/channel_state.cpp index a1a8110510..57d7dc8e57 100644 --- a/src/video_core/control/channel_state.cpp +++ b/src/video_core/control/channel_state.cpp @@ -13,7 +13,9 @@ #include "video_core/engines/maxwell_3d.h" #include "video_core/engines/maxwell_dma.h" #include "video_core/engines/puller.h" +#include "video_core/macro.h" #include "video_core/memory_manager.h" +#include "video_core/nx2.h" namespace Tegra::Control { namespace { @@ -25,6 +27,17 @@ constexpr u32 NvkM2mfSubchannel = 2; constexpr u32 Nvk2DSubchannel = 3; constexpr u32 NvkCopySubchannel = 4; +// Switch 2 channels of a process share one TSG, whose graphics context holds the macro RAM. +std::shared_ptr<NativeMacroContext> Nx2MacroContext() { + static const auto context = [] { + auto created = std::make_shared<NativeMacroContext>(); + created->engine.native_turing = true; + created->compute_engine.native_turing = true; + return created; + }(); + return context; +} + void BindNvkDefaultSubchannels(ChannelState::Payload& payload) { auto& dma_pusher = payload.dma_pusher; dma_pusher.BindSubchannel(&payload.maxwell_3d, Nvk3DSubchannel, Engines::EngineTypes::Maxwell3D); @@ -60,6 +73,16 @@ void ChannelState::Init(Core::System& system, u64 program_id_) { program_id = program_id_; payload.emplace(system, *memory_manager, *this); BindNvkDefaultSubchannels(*payload); + if (g_nx2_ampere) { + // Switch 2 command streams use Ampere classes with Turing+ (MME64) macros, and its + // block-linear surfaces the console's GOB swizzle. + memory_manager->EnableNativeTiling(); + const auto macros = Nx2MacroContext(); + payload->maxwell_3d.SetNativeMacroContext(macros); + payload->maxwell_3d.SetNativeTuringMacroMode(true); + payload->kepler_compute.SetNativeMacroMode(¯os->engine, macros->positions, + macros.get()); + } initialized = true; } diff --git a/src/video_core/dirty_flags.cpp b/src/video_core/dirty_flags.cpp index c2ecc12f56..e6856aaca3 100644 --- a/src/video_core/dirty_flags.cpp +++ b/src/video_core/dirty_flags.cpp @@ -22,6 +22,7 @@ void SetupDirtyVertexBuffers(Maxwell3D::DirtyState::Tables& tables) { OFF(vertex_stream_limits) + i * NUM(vertex_stream_limits[0]); FillBlock(tables, array_offset, num_array, VertexBuffer0 + i, VertexBuffers); + FillBlock(tables, 0x600 / sizeof(u32) + i * 2, 2, VertexBuffer0 + i, VertexBuffers); FillBlock(tables, limit_offset, NUM(vertex_stream_limits), VertexBuffer0 + i, VertexBuffers); } diff --git a/src/video_core/engines/engine_upload.cpp b/src/video_core/engines/engine_upload.cpp index ba21eec4c4..0ffdae8036 100644 --- a/src/video_core/engines/engine_upload.cpp +++ b/src/video_core/engines/engine_upload.cpp @@ -27,6 +27,7 @@ void State::BindRasterizer(VideoCore::RasterizerInterface* rasterizer_) { void State::ProcessExec(const bool is_linear_) { write_offset = 0; + flushed_offset = 0; copy_size = regs.line_length_in * regs.line_count; inner_buffer.resize_destructive(copy_size); is_linear = is_linear_; @@ -39,10 +40,29 @@ void State::ProcessData(const u32 data, const bool is_last_call) { if (!is_last_call) { return; } + if (native_partial && is_linear && regs.line_count == 1) { + const u32 count = write_offset - flushed_offset; + if (count) { + rasterizer->AccelerateInlineToMemory(regs.dest.Address() + flushed_offset, count, + std::span<const u8>{inner_buffer.data() + flushed_offset, count}); + flushed_offset = write_offset; + } + return; + } ProcessData(inner_buffer); } void State::ProcessData(const u32* data, size_t num_data) { + if (native_partial && is_linear && regs.line_count == 1) { + const auto count = std::min<size_t>(num_data * sizeof(u32), copy_size - write_offset); + if (count) { + rasterizer->AccelerateInlineToMemory(regs.dest.Address() + write_offset, count, + std::span<const u8>{reinterpret_cast<const u8*>(data), count}); + write_offset += static_cast<u32>(count); + flushed_offset = write_offset; + } + return; + } std::span<const u8> read_buffer(reinterpret_cast<const u8*>(data), num_data * sizeof(u32)); ProcessData(read_buffer); } diff --git a/src/video_core/engines/engine_upload.h b/src/video_core/engines/engine_upload.h index 21bf8aeb46..55a82f938b 100644 --- a/src/video_core/engines/engine_upload.h +++ b/src/video_core/engines/engine_upload.h @@ -62,6 +62,7 @@ public: explicit State(MemoryManager& memory_manager_, Registers& regs_); ~State(); + void SetNativePartialTransfers(bool enabled) { native_partial = enabled; } void ProcessExec(bool is_linear_); void ProcessData(u32 data, bool is_last_call); void ProcessData(const u32* data, size_t num_data); @@ -80,6 +81,8 @@ public: private: void ProcessData(std::span<const u8> read_buffer); + bool native_partial{}; + u32 flushed_offset{}; u32 write_offset = 0; u32 copy_size = 0; Common::ScratchBuffer<u8> inner_buffer; diff --git a/src/video_core/engines/fermi_2d.cpp b/src/video_core/engines/fermi_2d.cpp index aeff5e156f..a9b3355476 100644 --- a/src/video_core/engines/fermi_2d.cpp +++ b/src/video_core/engines/fermi_2d.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include <cstdlib> #include "common/assert.h" #include "common/logging.h" #include "video_core/engines/fermi_2d.h" @@ -59,6 +60,14 @@ void Fermi2D::ConsumeSinkImpl(Core::System& system) { } void Fermi2D::Blit() { + if (static const bool trace = std::getenv("EDEN_NX2_RT_TRACE") != nullptr; trace) { + static u32 traced{}; + if (traced++ < 64) { + LOG_INFO(HW_GPU, "nx2 2D blit {:#x} {}x{} -> {:#x} {}x{}", regs.src.Address(), + regs.src.width, regs.src.height, regs.dst.Address(), regs.dst.width, + regs.dst.height); + } + } LOG_DEBUG(HW_GPU, "called. source address={:#x}, destination address={:#x}", regs.src.Address(), regs.dst.Address()); UNIMPLEMENTED_IF_MSG(regs.operation != Operation::SrcCopy, "Operation is not copy"); diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 166bbd6d53..7710ea7dd2 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -5,12 +5,16 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include <bitset> +#include <cstdlib> +#include <fstream> #include "common/assert.h" #include "common/logging.h" +#include "shader_recompiler/exception.h" #include "core/core.h" #include "video_core/engines/kepler_compute.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/memory_manager.h" +#include "video_core/native_query_trace.h" #include "video_core/rasterizer_interface.h" #include "video_core/textures/decoders.h" @@ -42,15 +46,175 @@ void KeplerCompute::ConsumeSinkImpl(Core::System& system) { method_sink.clear(); } +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; + native_macro = native_compute_uploaded ? &native_context->compute_engine : &native_context->engine; + native_macro_positions = native_compute_uploaded ? std::span<u32>{native_context->compute_positions} : std::span<u32>{native_context->positions}; + } + if (method >= native_registers.size()) { + LOG_ERROR(HW_GPU, "Native compute method out of range: {:#x}", method); + return true; + } + if (method >= 0xe00) { + if (!native_executing_macro) { + if (method & 1) { LOG_ERROR(HW_GPU, "Native compute macro starts on argument method"); return true; } + native_executing_macro = method; + } + native_parameters.push_back(argument); + if (!is_last_call) return true; + const u32 entry = ((*native_executing_macro - 0xe00) / 2) % native_macro_positions.size(); + native_executing_macro.reset(); + 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 (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 (!result.success) { + native_macro->native_execution_failed = true; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-compute-mme.bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(native_macro->native_code.data()), sizeof(native_macro->native_code)); + } + LOG_ERROR(HW_GPU, "Native MME failed: compute start={:#x}, pc={:#x}, consumed={}/{}, error={}, parameters={:08x}", + native_macro_positions[entry], result.pc, result.consumed, native_parameters.size(), result.error, + fmt::join(native_parameters, ",")); + } + native_parameters.clear(); + return true; + } + native_registers[method] = argument; + if (method == 0x5a || (method >= 0xc6 && method <= 0x108)) { + VideoCommon::TraceNativeQuery(12, memory_manager.GetID(), method, 0, 0, argument, 0); + } + if (method == 0x107) { + const GPUVAddr address = ((GPUVAddr{native_registers[0xc6]} << 32) | + native_registers[0xc7]) << 8; + 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={}/{}/{}", + address, program, control, native_registers[0xc8 + 12], + native_registers[0xc8 + 13] & 0xffff, native_registers[0xc8 + 14] & 0xffff); + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-inline-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(native_registers.data() + 0xc8), 256); + std::ofstream registers("native-inline-compute-registers.bin", std::ios::binary); + registers.write(reinterpret_cast<const char*>(native_registers.data()), sizeof(native_registers)); + if (memory_manager.IsFullyMappedRange(program, 0x10000)) { + std::vector<u8> code(0x10000); + memory_manager.ReadBlockUnsafe(program, code.data(), code.size()); + std::ofstream shader("native-inline-compute-shader.bin", std::ios::binary); + shader.write(reinterpret_cast<const char*>(code.data()), code.size()); + } + } + if (!memory_manager.IsFullyMappedRange(address, 256)) { + native_macro->native_execution_failed = true; + LOG_ERROR(HW_GPU, "Unmapped native inline QMD destination"); + return true; + } + rasterizer->AccelerateInlineToMemory(address, 256, + std::span<const u8>{reinterpret_cast<const u8*>(native_registers.data() + 0xc8), 256}); + if ((control & (1U << 24)) || !(native_registers[0xcc] & (1U << 12))) { + const u32 previous = native_registers[0xad]; + native_registers[0xad] = static_cast<u32>(address >> 8); + ProcessNativeMethod(system, 0xb0, 3, true); + native_registers[0xad] = previous; + } + return true; + } + switch (method) { + case 0x45: native_macro->ClearCode(argument); return true; + case 0x46: native_macro->AddCode(0, argument); return true; + case 0x47: native_bind_pointer = argument; return true; + case 0x48: + if (native_bind_pointer < native_macro_positions.size()) native_macro_positions[native_bind_pointer++] = argument; + else LOG_ERROR(HW_GPU, "Native compute macro bind out of bounds"); + return true; + case 0xae: case 0xaf: case 0xb0: { + const GPUVAddr qmd_address = GPUVAddr{native_registers[0xad]} << 8; + VideoCommon::TraceNativeQuery(9, memory_manager.GetID(), qmd_address, + memory_manager.GpuToCpuAddress(qmd_address).value_or(0), method, argument, 0); + if (method == 0xae || (method == 0xaf && !(argument & 2)) || + (method == 0xb0 && ((argument & 15) == 0 || (argument & 15) == 1))) return true; + if ((argument & 15) == 12 && std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + static u32 captured{}; + if (captured++ == 0 && memory_manager.IsFullyMappedRange(qmd_address, 256)) { + std::array<u32, 64> descriptor{}; + memory_manager.ReadBlockUnsafe(qmd_address, descriptor.data(), sizeof(descriptor)); + std::ofstream dump("native-increment-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(descriptor.data()), sizeof(descriptor)); + LOG_WARNING(HW_GPU, "Native increment dependency descriptor: address={:#x}, words={:08x}", + qmd_address, fmt::join(descriptor, ",")); + } + } + const auto error = native_qmd.Signal(qmd_address, argument, + [&](u64 address, NativeQmdScheduler::Descriptor& descriptor) { + if (!memory_manager.IsFullyMappedRange(address, sizeof(descriptor))) return false; + memory_manager.ReadBlockUnsafe(address, descriptor.data(), sizeof(descriptor)); + for (u32 slot = 0; slot < 2; ++slot) { + const u32 flags = descriptor[25 + slot * 4]; + if (!(flags & (1U << 23))) continue; + const GPUVAddr destination = (GPUVAddr{flags & 0xff} << 32) | descriptor[24 + slot * 4]; + VideoCommon::TraceNativeQuery(7 + slot, memory_manager.GetID(), destination, + memory_manager.GpuToCpuAddress(destination).value_or(0), + descriptor[4], descriptor[26 + slot * 4], flags); + } + return true; + }, + [&](u64 address, u32 payload, u32 flags) { + if (!memory_manager.IsFullyMappedRange(address, sizeof(payload))) return false; + if (flags & (1U << 26)) { + rasterizer->SignalFence([this, address, payload] { + const u32 value = memory_manager.Read<u32>(address) + payload; + memory_manager.Write<u32>(address, value); + VideoCommon::TraceNativeQuery(1, memory_manager.GetID(), address, + memory_manager.GpuToCpuAddress(address).value_or(0), 0, value, 1U << 26); + }); + } else { + rasterizer->Query(address, VideoCommon::QueryType::Payload, + VideoCommon::QueryPropertiesFlags::IsAFence, payload, 0); + } + return true; + }, [&](const auto& descriptor) { return DispatchNativeCompute(descriptor); }); + if (error.empty()) { + LOG_DEBUG(HW_GPU, "Native QMD control action={:#x}, QMD={:#x}", argument, qmd_address); + return true; + } + native_macro->native_execution_failed = true; + LOG_ERROR(HW_GPU, "Native QMD failed: {}", error); + std::array<u32, 128> qmd{}; + if (memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { + memory_manager.ReadBlockUnsafe(qmd_address, qmd.data(), sizeof(qmd)); + } + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && memory_manager.IsFullyMappedRange(qmd_address, sizeof(qmd))) { + std::ofstream dump("native-qmd.bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(qmd.data()), sizeof(qmd)); + std::ofstream macros("native-dispatch-mme.bin", std::ios::binary); + macros.write(reinterpret_cast<const char*>(native_macro->native_code.data()), sizeof(native_macro->native_code)); + std::ofstream registers("native-dispatch-registers.bin", std::ios::binary); + registers.write(reinterpret_cast<const char*>(native_registers.data()), sizeof(native_registers)); + } + LOG_ERROR(HW_GPU, "Native compute dispatch unsupported: method={:#x}, value={:#x}, QMD={:#x}, words={:08x}", method, argument, + qmd_address, fmt::join(qmd, ",")); + return true; + } + default: + // The inline upload register block is shared with the older class. + if (method < Regs::NUM_REGS) regs.reg_array[method] = argument; + return method != KEPLER_COMPUTE_REG_INDEX(exec_upload) && method != KEPLER_COMPUTE_REG_INDEX(data_upload); + } +} + void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) { + if (native_mode && ProcessNativeMethod(system, method, method_argument, is_last_call)) return; 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); - } + LOG_ERROR(HW_GPU, "Compute register out of range: {:#x}", method); return; } @@ -96,6 +260,10 @@ void KeplerCompute::CallMethod(Core::System& system, u32 method, u32 method_argu } void KeplerCompute::CallMultiMethod(Core::System& system, u32 method, const u32* base_start, u32 amount, u32 methods_pending) { + if (native_mode) { + for (u32 i = 0; i < amount; ++i) CallMethod(system, method, base_start[i], methods_pending - i <= 1); + return; + } switch (method) { case KEPLER_COMPUTE_REG_INDEX(data_upload): upload_address = current_dma_segment; @@ -111,6 +279,42 @@ void KeplerCompute::CallMultiMethod(Core::System& system, u32 method, const u32* } } +bool KeplerCompute::DispatchNativeCompute(const NativeQmdScheduler::Descriptor& qmd) { + launch_description = {}; + launch_description.grid_dim_x.Assign(qmd[12]); + launch_description.grid_dim_y.Assign(qmd[13] & 0xffff); + launch_description.grid_dim_z.Assign(qmd[14] & 0xffff); + launch_description.shared_alloc.Assign(qmd[17] & 0x3ffff); + launch_description.block_dim_x.Assign(qmd[18] >> 16); + launch_description.block_dim_y.Assign(qmd[19] & 0xffff); + launch_description.block_dim_z.Assign(qmd[19] >> 16); + launch_description.const_buffer_enable_mask.Assign(qmd[20] & 0xff); + launch_description.local_pos_alloc.Assign(qmd[23] & 0xffffff); + launch_description.local_crs_alloc.Assign(qmd[50] & 0xffffff); + regs.code_loc.address_low = qmd[48]; + regs.code_loc.address_high = qmd[49] & 0x1ffff; + for (u32 slot = 0; slot < NumConstBuffers; ++slot) { + const u32 high = qmd[33 + slot * 2]; + if (high & 0x1ff00) { + LOG_ERROR(HW_GPU, "Native compute constant address exceeds supported GPU width"); + return false; + } + auto& buffer = launch_description.const_buffer_config[slot]; + buffer.address_low = qmd[32 + slot * 2]; + buffer.address_high.Assign(high & 0xff); + buffer.size.Assign(((high >> 19) & 0x1fff) << 4); + } + if (!launch_description.block_dim_x || !launch_description.block_dim_y || !launch_description.block_dim_z) + return false; + try { + rasterizer->DispatchCompute(); + } catch (const Shader::Exception& error) { + LOG_ERROR(HW_GPU, "Native compute shader rejected: {}", error.what()); + return false; + } + return true; +} + void KeplerCompute::ProcessLaunch() { const GPUVAddr launch_desc_loc = regs.launch_desc_loc.Address(); memory_manager.ReadBlockUnsafe(launch_desc_loc, &launch_description, diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index 078d3d45be..99cf74d1e8 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -15,6 +15,8 @@ #include "common/common_types.h" #include "video_core/engines/engine_interface.h" #include "video_core/engines/engine_upload.h" +#include "video_core/macro.h" +#include "video_core/native_qmd.h" #include "video_core/textures/texture.h" namespace Core { @@ -51,6 +53,22 @@ public: /// Binds a rasterizer to this engine. void BindRasterizer(VideoCore::RasterizerInterface* rasterizer); + void SetNativeMacroMode(MacroEngine* shared_macro, std::span<u32> positions, NativeMacroContext* context = nullptr) { + if (native_context != context) native_compute_uploaded = false; + native_context = context; + native_mode = shared_macro != nullptr; + upload_state.SetNativePartialTransfers(native_mode); + native_macro = shared_macro; + native_macro_positions = positions; + if (native_mode) execution_mask.set(); + else { + execution_mask.reset(); + execution_mask[KEPLER_COMPUTE_REG_INDEX(exec_upload)] = true; + execution_mask[KEPLER_COMPUTE_REG_INDEX(data_upload)] = true; + execution_mask[KEPLER_COMPUTE_REG_INDEX(launch)] = true; + } + } + static constexpr std::size_t NumConstBuffers = 8; struct Regs { @@ -211,7 +229,22 @@ public: return indirect_compute; } + bool IsNativeMode() const { return native_mode; } + private: + bool DispatchNativeCompute(const NativeQmdScheduler::Descriptor& descriptor); + bool ProcessNativeMethod(Core::System& system, u32 method, u32 argument, bool is_last_call); + NativeQmdScheduler native_qmd; + bool native_mode{}; + NativeMacroContext* native_context{}; + bool native_compute_uploaded{}; + MacroEngine* native_macro{}; + std::array<u32, 0x1000> native_registers{}; + std::span<u32> native_macro_positions; + std::vector<u32> native_parameters; + std::optional<u32> native_executing_macro; + u32 native_bind_pointer{}; + void ProcessLaunch(); void ConsumeSinkImpl(Core::System& system) override; diff --git a/src/video_core/engines/maxwell_3d.cpp b/src/video_core/engines/maxwell_3d.cpp index 1a43133296..f86b8e9351 100644 --- a/src/video_core/engines/maxwell_3d.cpp +++ b/src/video_core/engines/maxwell_3d.cpp @@ -5,6 +5,7 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include <cstring> +#include <limits> #include <optional> #include "common/assert.h" #include "common/bit_util.h" @@ -12,7 +13,6 @@ #include "common/settings.h" #include "core/core.h" #include "core/core_timing.h" -#include "video_core/nx2.h" #include "video_core/dirty_flags.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/engines/maxwell_3d.h" @@ -254,7 +254,10 @@ u32 Maxwell3D::GetMaxCurrentVertices() { const auto& limit = regs.vertex_stream_limits[index]; const GPUVAddr gpu_addr_begin = array.Address(); const GPUVAddr gpu_addr_end = limit.Address() + 1; - const u32 address_size = static_cast<u32>(gpu_addr_end - gpu_addr_begin); + const u64 stream_size = GetNativeMacroEngine().native_turing + ? regs.NativeVertexStreamSize(index) + : gpu_addr_end - gpu_addr_begin; + const u32 address_size = static_cast<u32>(std::min<u64>(stream_size, std::numeric_limits<u32>::max())); num_vertices = (std::max)( num_vertices, address_size / (std::max)(attribute.SizeInBytes(), array.stride.Value())); break; @@ -306,6 +309,9 @@ void Maxwell3D::ProcessDirtyRegisters(u32 method, u32 argument) { } void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argument, bool is_last_call) { + if (GetNativeMacroEngine().native_turing && method == MAXWELL3D_REG_INDEX(clear_surface)) { + LOG_DEBUG(HW_GPU, "Native GPU surface clear: value={:#x}", argument); + } switch (method) { case MAXWELL3D_REG_INDEX(wait_for_idle): return rasterizer->WaitForIdle(); @@ -313,9 +319,9 @@ void Maxwell3D::ProcessMethodCall(u32 method, u32 argument, u32 nonshadow_argume shadow_state.shadow_ram_control = static_cast<Regs::ShadowRamControl>(nonshadow_argument); return; case MAXWELL3D_REG_INDEX(load_mme.instruction_ptr): - return macro_engine.ClearCode(regs.load_mme.instruction_ptr); + return GetNativeMacroEngine().ClearCode(regs.load_mme.instruction_ptr); case MAXWELL3D_REG_INDEX(load_mme.instruction): - return macro_engine.AddCode(regs.load_mme.instruction_ptr, argument); + return GetNativeMacroEngine().AddCode(regs.load_mme.instruction_ptr, argument); case MAXWELL3D_REG_INDEX(load_mme.start_address): return ProcessMacroBind(argument); case MAXWELL3D_REG_INDEX(falcon[4]): @@ -379,16 +385,10 @@ void Maxwell3D::CallMacroMethod(Core::System& system, u32 method, const std::vec // Lookup the macro offset const u32 entry = - ((method - MacroRegistersStart) >> 1) % static_cast<u32>(macro_positions.size()); - - // Switch 2 uploads Turing+ (MME64) macros, which Eden's Fermi-style macro engine can't run. - if (g_nx2_ampere) { - LOG_DEBUG(HW_GPU, "Switch 2 macro {} with {} parameters skipped", entry, parameters.size()); - return; - } + ((method - MacroRegistersStart) >> 1) % static_cast<u32>(GetNativeMacroPositions().size()); // Execute the current macro. - macro_engine.Execute(system, *this, macro_positions[entry], parameters); + GetNativeMacroEngine().Execute(system, *this, GetNativeMacroPositions()[entry], parameters); draw_manager.DrawDeferred(*this); } @@ -460,11 +460,11 @@ void Maxwell3D::CallMultiMethod(Core::System& system, u32 method, const u32* bas } void Maxwell3D::ProcessMacroUpload(u32 data) { - macro_engine.AddCode(regs.load_mme.instruction_ptr++, data); + GetNativeMacroEngine().AddCode(regs.load_mme.instruction_ptr++, data); } void Maxwell3D::ProcessMacroBind(u32 data) { - macro_positions[regs.load_mme.start_address_ptr++] = data; + GetNativeMacroPositions()[regs.load_mme.start_address_ptr++] = data; } void Maxwell3D::ProcessFirmwareCall4() { @@ -593,6 +593,9 @@ void Maxwell3D::ProcessCounterReset() { } void Maxwell3D::ProcessSyncPoint() { + if (GetNativeMacroEngine().native_turing) { + LOG_DEBUG(HW_GPU, "Native 3D syncpoint increment: id={}, raw={:#x}", regs.sync_info.sync_point.Value(), regs.reg_array[MAXWELL3D_REG_INDEX(sync_info)]); + } const u32 sync_point = regs.sync_info.sync_point.Value(); [[maybe_unused]] const u32 cache_flush = regs.sync_info.clean_l2.Value(); rasterizer->SignalSyncPoint(sync_point); diff --git a/src/video_core/engines/maxwell_3d.h b/src/video_core/engines/maxwell_3d.h index 0cab2ae986..8cbc4f4274 100644 --- a/src/video_core/engines/maxwell_3d.h +++ b/src/video_core/engines/maxwell_3d.h @@ -2605,6 +2605,12 @@ public: u32 stop_shader_mask; }; + // Native C597/C797 use byte counts at 0x600 rather than legacy end addresses. + u64 NativeVertexStreamSize(std::size_t index) const { + const std::size_t method = 0x600 / sizeof(u32) + index * 2; + return (u64{reg_array[method] & 0xffU} << 32) | reg_array[method + 1]; + } + // clang-format off union { struct { @@ -3129,6 +3135,17 @@ public: /// Reads a register value located at the input method address u32 GetRegisterValue(u32 method) const; + void SetNativeTuringMacroMode(bool enabled) { + GetNativeMacroEngine().native_turing = enabled; + upload_state.SetNativePartialTransfers(enabled); + } + void SetNativeMacroContext(std::shared_ptr<NativeMacroContext> context) { + native_macro_context = std::move(context); + } + MacroEngine& GetNativeMacroEngine() { return native_macro_context ? native_macro_context->engine : macro_engine; } + std::span<u32> GetNativeMacroPositions() { return native_macro_context ? std::span<u32>{native_macro_context->positions} : std::span<u32>{macro_positions}; } + + /// Write the value to the register identified by method. void CallMethod(Core::System& system, u32 method, u32 method_argument, bool is_last_call) override; @@ -3264,6 +3281,7 @@ private: std::vector<u32> macro_params; /// Interpreter for the macro codes uploaded to the GPU. + std::shared_ptr<NativeMacroContext> native_macro_context; MacroEngine macro_engine; Upload::State upload_state; diff --git a/src/video_core/engines/maxwell_dma.cpp b/src/video_core/engines/maxwell_dma.cpp index 5ab953e510..5b6efabca0 100644 --- a/src/video_core/engines/maxwell_dma.cpp +++ b/src/video_core/engines/maxwell_dma.cpp @@ -4,6 +4,7 @@ // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include <cstdlib> #include "common/algorithm.h" #include "common/assert.h" #include "common/logging.h" @@ -63,6 +64,14 @@ void MaxwellDMA::CallMultiMethod(Core::System& system, u32 method, const u32* ba void MaxwellDMA::Launch() { LOG_TRACE(Render_OpenGL, "DMA copy {:#x} -> {:#x}", static_cast<GPUVAddr>(regs.offset_in), GPUVAddr(regs.offset_out)); + if (static const bool trace = std::getenv("EDEN_NX2_RT_TRACE") != nullptr; trace) { + static u32 traced{}; + if (traced++ < 64) { + LOG_INFO(HW_GPU, "nx2 DMA {:#x} -> {:#x} lines={} length={} launch={:#x}", + static_cast<GPUVAddr>(regs.offset_in), GPUVAddr(regs.offset_out), + regs.line_count, regs.line_length_in, regs.reg_array[0xc0]); + } + } // TODO(Subv): Perform more research and implement all features of this engine. const LaunchDMA& launch = regs.launch_dma; @@ -223,7 +232,7 @@ void MaxwellDMA::CopyBlockLinearToPitch() { UnswizzleSubrect(tmp_write_buffer, tmp_read_buffer, bytes_per_pixel, width, height, depth, x_offset, src_params.origin.y, x_elements, regs.line_count, block_height, - block_depth, dst_operand.pitch); + block_depth, dst_operand.pitch, memory_manager.UsesNativeTiling()); } void MaxwellDMA::CopyPitchToBlockLinear() { @@ -288,7 +297,7 @@ void MaxwellDMA::CopyPitchToBlockLinear() { // If the input is linear and the output is tiled, swizzle the input and copy it over. SwizzleSubrect(tmp_write_buffer, tmp_read_buffer, bytes_per_pixel, width, height, depth, x_offset, dst_params.origin.y, x_elements, regs.line_count, block_height, - block_depth, regs.pitch_in); + block_depth, regs.pitch_in, memory_manager.UsesNativeTiling()); } void MaxwellDMA::CopyBlockLinearToBlockLinear() { @@ -341,11 +350,11 @@ void MaxwellDMA::CopyBlockLinearToBlockLinear() { UnswizzleSubrect(intermediate_buffer, tmp_read_buffer, bytes_per_pixel, src_width, src.height, src.depth, src_x_offset, src.origin.y, x_elements, regs.line_count, - src.block_size.height, src.block_size.depth, pitch); + src.block_size.height, src.block_size.depth, pitch, memory_manager.UsesNativeTiling()); SwizzleSubrect(tmp_write_buffer, intermediate_buffer, bytes_per_pixel, dst_width, dst.height, dst.depth, dst_x_offset, dst.origin.y, x_elements, regs.line_count, - dst.block_size.height, dst.block_size.depth, pitch); + dst.block_size.height, dst.block_size.depth, pitch, memory_manager.UsesNativeTiling()); } void MaxwellDMA::ReleaseSemaphore() { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_2d.comp b/src/video_core/host_shaders/block_linear_unswizzle_2d.comp index 2a3446ca66..bbe73acd1a 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_2d.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_2d.comp @@ -39,6 +39,7 @@ UNIFORM(4) uint block_size; UNIFORM(5) uint x_shift; UNIFORM(6) uint block_height; UNIFORM(7) uint block_height_mask; +UNIFORM(8) uint native_tiling; END_PUSH_CONSTANTS #if HAS_EXTENDED_TYPES @@ -77,7 +78,12 @@ uint SwizzleTable(uint pos) { uint SwizzleOffset(uvec2 pos) { pos = pos & SWIZZLE_MASK; - return SwizzleTable(pos.y * 64 + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_3d.comp b/src/video_core/host_shaders/block_linear_unswizzle_3d.comp index 34e4f93307..8dfaf0cd1a 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_3d.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_3d.comp @@ -41,6 +41,7 @@ UNIFORM(6) uint block_height; UNIFORM(7) uint block_height_mask; UNIFORM(8) uint block_depth; UNIFORM(9) uint block_depth_mask; +UNIFORM(10) uint native_tiling; END_PUSH_CONSTANTS #if HAS_EXTENDED_TYPES @@ -79,7 +80,12 @@ uint SwizzleTable(uint pos) { uint SwizzleOffset(uvec2 pos) { pos = pos & SWIZZLE_MASK; - return SwizzleTable(pos.y * 64 + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp index e084837f25..28b1dcf599 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp @@ -42,7 +42,7 @@ layout(push_constant) uniform PushConstants { uint block_depth; // Offset 48 uint block_depth_mask; // Offset 52 - int _pad; // Offset 56 + uint native_tiling; // Offset 56 ivec3 destination; // Offset 60 } pc; @@ -59,6 +59,7 @@ BEGIN_PUSH_CONSTANTS UNIFORM(8) uint block_depth; UNIFORM(9) uint block_depth_mask; UNIFORM(10) uvec3 blocks_dim; + UNIFORM(11) uint native_tiling; END_PUSH_CONSTANTS #define pc // Map pc prefix to nothing for OpenGL compatibility #endif @@ -103,7 +104,12 @@ uint SwizzleTable(uint pos) { // --- Helpers --- uint SwizzleOffset(uvec2 pos) { pos &= SWIZZLE_MASK; - return SwizzleTable(pos.y * 64u + pos.x); + uint offset = SwizzleTable(pos.y * 64 + pos.x); + if (pc.native_tiling != 0) { + uint difference = ((offset >> 5) ^ (offset >> 6)) & 1; + offset ^= (difference << 5) | (difference << 6); + } + return offset; } uvec4 ReadTexel(uint offset) { diff --git a/src/video_core/macro.cpp b/src/video_core/macro.cpp index 86f528af15..a8acd81682 100644 --- a/src/video_core/macro.cpp +++ b/src/video_core/macro.cpp @@ -4,6 +4,8 @@ // SPDX-FileCopyrightText: Copyright 2020 yuzu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +#include <cstdlib> +#include <unordered_set> #include <cstring> #include <fstream> #include <optional> @@ -1339,6 +1341,34 @@ static void Dump(u64 hash, std::span<const u32> code, bool decompiled = false) { void MacroEngine::Execute(Core::System& system, Engines::Maxwell3D& maxwell3d, u32 method, std::span<const u32> parameters) { + if (native_turing) { + maxwell3d.RefreshParameters(); + const auto context = reinterpret_cast<std::uintptr_t>(this); + static std::unordered_set<std::uintptr_t> captured_contexts; + if (std::getenv("EDEN_NATIVE_FRAME_CAPTURE") && captured_contexts.insert(context).second) { + std::ofstream dump("native-graphics-mme-" + std::to_string(context) + ".bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(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, ",")); + const auto result = native_interpreter.Execute(native_code, method, parameters, + [&](u32 address) { return address < Maxwell3D::Regs::NUM_REGS ? maxwell3d.GetRegisterValue(address) : 0U; }, + [&](u32 address, u32 value) { + if (address >= 0xe00) return false; + maxwell3d.CallMethod(system, address, value, true); + return true; + }); + if (!result.success) { + native_execution_failed = true; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + std::ofstream dump("native-mme-" + std::to_string(method) + ".bin", std::ios::binary); + dump.write(reinterpret_cast<const char*>(native_code.data()), sizeof(native_code)); + } + LOG_ERROR(HW_GPU, "Native MME failed: start={:#x}, pc={:#x}, consumed={}/{}, error={}", + method, result.pc, result.consumed, parameters.size(), result.error); + } + return; + } const auto execute_variant = [&system, &maxwell3d, ¶meters, method](AnyCachedMacro& cached) { if (std::holds_alternative<MacroInterpreterImpl>(cached) || diff --git a/src/video_core/macro.h b/src/video_core/macro.h index 90519a5429..52dfe7c82c 100644 --- a/src/video_core/macro.h +++ b/src/video_core/macro.h @@ -7,6 +7,7 @@ #pragma once #include <memory> +#include "video_core/macro_turing.h" #include <span> #include <variant> #include <vector> @@ -223,10 +224,15 @@ struct MacroEngine { MacroEngine(bool is_interpreted_) noexcept : is_interpreted{is_interpreted_} {} // Store the uploaded macro code to compile them when they're called. inline void AddCode(u32 method, u32 data) noexcept { + if (native_turing) { + if (native_write_offset < native_code.size()) native_code[native_write_offset++] = data; + return; + } uploaded_macro_code[method].push_back(data); } // Clear the code associated with a method. inline void ClearCode(u32 method) noexcept { + if (native_turing) { native_write_offset = method < 0x1000 ? method * 3 : static_cast<u32>(native_code.size()); return; } macro_cache.erase(method); uploaded_macro_code.erase(method); } @@ -240,6 +246,19 @@ struct MacroEngine { ::Common::unordered_map<u32, CacheInfo> macro_cache; ::Common::unordered_map<u32, std::vector<u32>> uploaded_macro_code; bool is_interpreted; + bool native_turing{}; + bool native_execution_failed{}; + u32 native_write_offset{}; + std::array<u32, 0x3000> native_code{}; + TuringMacroInterpreter native_interpreter; + +}; + +struct NativeMacroContext { + MacroEngine engine{true}; + std::array<u32, 0x80> positions{}; + MacroEngine compute_engine{true}; + std::array<u32, 0x80> compute_positions{}; }; } // namespace Tegra diff --git a/src/video_core/macro_turing.h b/src/video_core/macro_turing.h new file mode 100644 index 0000000000..fe82ce4a4f --- /dev/null +++ b/src/video_core/macro_turing.h @@ -0,0 +1,213 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +// Instruction semantics adapted from Mesa mme_tu104_sim.c and mme_tu104.c: +// Copyright © 2022 Collabora Ltd. SPDX-License-Identifier: MIT +#pragma once + +#include <algorithm> +#include <array> +#include <bit> +#include <cstdint> +#include <span> +#include <string_view> + +namespace Tegra { + +// TU104/GA10x MME has two ALUs and two outputs per 96-bit instruction. +class TuringMacroInterpreter { +public: + using Word = std::uint32_t; + struct Result { + bool success; + Word pc; + std::size_t consumed; + std::string_view error; + }; + + template <typename Read, typename Emit> + Result Execute(std::span<const Word> code, Word start, + std::span<const Word> parameters, Read read, Emit emit) { + std::array<Word, 24> regs{}; + Word initialized{}, carry{}, method{}, increment{}; + Word pc = start, loop_count{}, loop_start{}, loop_end{}; + std::size_t consumed{}; + bool end_next{}, stopped{}; + std::string_view error; + const auto fail = [&](std::string_view message) { if (error.empty()) error = message; }; + for (Word steps = 0; steps < 1000000; ++steps) { + if (static_cast<std::size_t>(pc) * 3 + 2 >= code.size()) { + fail("instruction outside uploaded RAM"); + break; + } + const std::array<Word, 3> bits{code[pc * 3 + 2], code[pc * 3 + 1], code[pc * 3]}; + const auto field = [&](Word low, Word count) { + std::uint64_t value = bits[low / 32]; + if (low / 32 < 2) value |= std::uint64_t{bits[low / 32 + 1]} << 32; + return Word((value >> (low % 32)) & ((std::uint64_t{1} << count) - 1)); + }; + const std::array<Word, 2> imm{field(30, 16), field(66, 16)}; + const std::array<Word, 2> op{field(10, 5), field(46, 5)}; + const std::array<Word, 2> dst{field(15, 5), field(51, 5)}; + const std::array<Word, 2> src_a{field(20, 5), field(56, 5)}; + const std::array<Word, 2> src_b{field(25, 5), field(61, 5)}; + const std::array<Word, 2> out_method{field(82, 3), field(89, 3)}; + const std::array<Word, 2> out_emit{field(85, 4), field(92, 4)}; + const Word pred_reg = field(5, 5), pred_mode = field(1, 4); + std::array<Word, 2> load{}, alu{}; + for (Word i = 0; i < 2; ++i) { + const Word reg = 28 + i, out = 3 + i; + const bool needed = pred_reg == reg || src_a[0] == reg || src_a[1] == reg || + src_b[0] == reg || src_b[1] == reg || out_method[0] == out || + out_method[1] == out || out_emit[0] == out || out_emit[1] == out; + if (needed) { + if (consumed >= parameters.size()) { fail("parameter FIFO underflow"); break; } + load[i] = parameters[consumed++]; + } + } + // Both ALUs read the register file before either destination is + // committed. Carry chaining remains explicit between the lanes. + const auto input_regs = regs; + const Word input_initialized = initialized; + const auto get = [&](Word reg, Word lane) -> Word { + if (reg < 24) { + if (!(input_initialized & (1U << reg))) fail("uninitialized register read"); + return input_regs[reg]; + } + switch (reg) { + case 24: return 0; + case 25: return Word(std::int32_t(std::int16_t(imm[lane]))); + case 26: return Word(std::int32_t(std::int16_t(imm[1 - lane]))); + case 27: return (imm[0] << 16) | imm[1]; + case 28: return load[0]; + case 29: return load[1]; + default: fail("invalid source register"); return 0; + } + }; + constexpr std::array<std::string_view, 16> predicates{ + "UUUU", "TTTT", "FFFF", "TTUU", "FFUU", "TFUU", "TUUU", "FUUU", + "UUTT", "UUTF", "UUTU", "UUFT", "UUFF", "UUFU", "UUUT", "UUUF"}; + Word mask = 15; + if (pred_mode) { + const char skipped = get(pred_reg, 0) ? 'T' : 'F'; + mask = 0; + for (Word i = 0; i < 4; ++i) + if (predicates[pred_mode][i] != skipped) mask |= 1U << i; + } + Word next_pc = pc + 1; + bool branch_cancels_end{}; + for (Word lane = 0; lane < 2 && error.empty(); ++lane) { + if (!(mask & (1U << lane))) continue; + const Word x = get(src_a[lane], lane), y = get(src_b[lane], lane); + Word value{}; + const auto condition = [&](Word comparison) { + switch (comparison) { + case 17: case 25: return std::int32_t(x) < std::int32_t(y); + case 18: case 26: return x < y; + case 19: case 27: return std::int32_t(x) <= std::int32_t(y); + case 20: case 28: return x <= y; + default: return x == y; + } + }; + switch (op[lane]) { + case 0: value = x + y; carry = value < x; break; + case 1: value = x + y + carry; break; + case 2: value = x - y; carry = value > x; break; + case 3: value = x - y - carry; break; + case 4: case 6: { + const std::uint64_t a = op[lane] == 4 ? std::uint64_t(std::int64_t(std::int32_t(x))) : x; + const std::uint64_t b = op[lane] == 4 ? std::uint64_t(std::int64_t(std::int32_t(y))) : y; + const auto product = a * b; + value = Word(product); carry = Word(product >> 32); break; + } + case 5: value = carry; break; + case 7: + // EXTENDED: x = 0x1000 is the load barrier (Mesa mme_tu104_load_barrier), + // nothing to do in a sequential interpreter. + if (x != 0x1000) fail("unsupported extended operation"); + break; + case 8: value = Word(std::countl_zero(x)); break; + case 9: value = x << (y & 31); break; + case 10: value = x >> (y & 31); break; + case 11: value = Word(std::int32_t(x) >> (y & 31)); break; + case 12: value = x & y; break; + case 13: value = ~(x & y); break; + case 14: value = x | y; break; + case 15: value = x ^ y; break; + case 16: { + const Word dest = (imm[lane] >> 10) & 63, count = (imm[lane] >> 5) & 31; + const Word source = imm[lane] & 31; + if (dest >= 32) { fail("invalid merge destination"); break; } + const Word bits_mask = (1U << count) - 1; + value = (x & ~(bits_mask << dest)) | (((y >> source) & bits_mask) << dest); + break; + } + case 17: case 18: case 19: case 20: case 21: + value = condition(op[lane]) ? ~0U : 0; break; + case 22: value = read((x + y) & 0xffff); break; + case 23: + if (loop_count || imm[lane] < 2) { fail("invalid loop"); break; } + // The native SDK emits zero-count loops around FIFO reads. + // Unlike Mesa's test simulator, these must skip the body. + if (x == 0) { next_pc = pc + imm[lane]; break; } + loop_count = x - 1; loop_start = pc; loop_end = pc + imm[lane] - 1; + break; + case 24: next_pc = pc + (imm[lane] & 0x7fff); break; + case 25: case 26: case 27: case 28: case 29: + if (condition(op[lane]) == bool(imm[lane] & 0x8000)) { + const auto offset = std::int32_t(imm[lane] << 19) >> 19; + if (offset == -4096) stopped = true; + else { + next_pc = Word(std::int32_t(pc) + offset); + // Native back edges use bit 14 to suppress the + // fallthrough termination when the branch is taken. + branch_cancels_end = (imm[lane] & 0x4000) != 0; + } + } + break; + case 30: + if (x >= dram.size()) { fail("data RAM read out of bounds"); break; } + value = dram[x]; break; + case 31: + if (x >= dram.size()) { fail("data RAM write out of bounds"); break; } + dram[x] = y; break; + } + alu[lane] = value; + if (dst[lane] < 24) { regs[dst[lane]] = value; initialized |= 1U << dst[lane]; } + else if (dst[lane] != 24) fail("invalid destination register"); + } + const auto output = [&](Word selector) -> Word { + switch (selector) { + case 1: return alu[0]; case 2: return alu[1]; + case 3: return load[0]; case 4: return load[1]; + case 5: return imm[0]; case 6: return imm[1]; + case 8: return imm[0] >> 12; case 9: return imm[1] >> 12; + case 10: return (imm[0] << 16) | imm[1]; + default: fail("invalid output selector"); return 0; + } + }; + for (Word lane = 0; lane < 2 && error.empty(); ++lane) { + if (!(mask & (4U << lane))) continue; + if (out_method[lane]) { + const Word address = output(out_method[lane]); + method = address & 0xfff; increment = (address >> 12) & 15; + } + if (out_emit[lane]) { + const Word value = output(out_emit[lane]); + if (error.empty() && !emit(method, value)) fail("unsupported emitted method"); + method += increment; + } + } + if (!error.empty()) break; + if (end_next || stopped) return {true, pc, consumed, {}}; + end_next = field(0, 1) != 0 && !branch_cancels_end; + if (loop_count && pc == loop_end) { --loop_count; next_pc = loop_start + 1; } + pc = next_pc; + } + if (error.empty()) error = "instruction limit exceeded"; + return {false, pc, consumed, error}; + } + + std::array<Word, 0xc00> dram{}; +}; + +} // namespace Tegra diff --git a/src/video_core/memory_manager.cpp b/src/video_core/memory_manager.cpp index 52e9db189d..9a9fdab7bc 100644 --- a/src/video_core/memory_manager.cpp +++ b/src/video_core/memory_manager.cpp @@ -22,6 +22,10 @@ #include "video_core/texture_cache/util.h" namespace Tegra { +bool MemoryManager::IsShuttingDown() const { + return system.IsShuttingDown(); +} + using Tegra::Memory::GuestMemoryFlags; std::atomic<size_t> MemoryManager::unique_identifier_generator{}; diff --git a/src/video_core/memory_manager.h b/src/video_core/memory_manager.h index 1aaf2c1086..105cb14e23 100644 --- a/src/video_core/memory_manager.h +++ b/src/video_core/memory_manager.h @@ -47,6 +47,15 @@ public: void SetNativeSyncpointShim(GPUVAddr address); static constexpr u32 NumShimSyncpoints = 192; // Tegra::Host1x::SyncpointManager + /// Switch 2 lays out block-linear surfaces with the console's GOB swizzle. + void EnableNativeTiling() { + native_tiling = true; + } + bool UsesNativeTiling() const { + return native_tiling; + } + bool IsShuttingDown() const; + static constexpr bool HAS_FLUSH_INVALIDATION = true; inline size_t GetID() const noexcept { @@ -225,6 +234,7 @@ private: boost::container::small_vector<std::pair<DAddr, std::size_t>, 32> page_stash{}; boost::container::small_vector<std::pair<DAddr, std::size_t>, 32> page_stash2{}; + bool native_tiling{}; mutable std::mutex guard; static constexpr size_t continuous_bits = 64; diff --git a/src/video_core/native_qmd.h b/src/video_core/native_qmd.h new file mode 100644 index 0000000000..c9ae3b8cc0 --- /dev/null +++ b/src/video_core/native_qmd.h @@ -0,0 +1,112 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once +#include <array> +#include <cstdint> +#include <string_view> +#include <type_traits> +#include <unordered_map> + +namespace Tegra { + +// Native QMD 3.0 dependency scheduling; nonempty grids require a dispatch backend. +class NativeQmdScheduler { +public: + using Word = std::uint32_t; + using Address = std::uint64_t; + using Descriptor = std::array<Word, 64>; + struct RejectShaderDispatch { bool operator()(const Descriptor&) const { return false; } }; + template <typename Read, typename Release> + std::string_view Signal(Address address, Word argument, Read read, Release release) { + return Signal(address, argument, read, release, RejectShaderDispatch{}); + } + template <typename Read, typename Release, typename Dispatch> + std::string_view Signal(Address address, Word argument, Read read, Release release, Dispatch dispatch) { + struct Pending { Address address; Word action; Word count; }; + std::array<Pending, 256> pending{}; + std::size_t count = 1; + pending[0] = {address, argument & 15, ((argument >> 14) & 0x3ff) + 1}; + for (std::size_t i = 0; i < count; ++i) { + const auto request = pending[i]; + if (states.size() >= 4096 && !states.contains(request.address)) return "QMD cache limit"; + auto& state = states[request.address]; + if (request.action == 3 || request.action == 2) { + if (request.action == 3 || !state.loaded) { + 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"; + if constexpr (std::is_same_v<Dispatch, RejectShaderDispatch>) { + if (qmd[12] && (qmd[13] & 0xffff) && (qmd[14] & 0xffff)) + return "native shader dispatch required"; + } + state.dependencies = std::int32_t((qmd[4] >> 14) & 0xffff) + state.deferred; + state.deferred = 0; + state.loaded = true; + state.completed = false; + } + state.scheduled = true; + } else if (request.action == 7 || request.action == 12) { + const auto delta = request.action == 7 ? -std::int32_t(request.count) : std::int32_t(request.count); + if (!state.loaded || state.completed) { state.deferred += delta; continue; } + state.dependencies += delta; + } else if (request.action == 0 || request.action == 1) { + continue; + } else { + return "unsupported PCAS action"; + } + if (!state.loaded || !state.scheduled || state.completed || state.dependencies > 0) continue; + if (state.dependencies < 0) return "QMD dependency underflow"; + // The SDK patches release payloads while a scheduled QMD is + // waiting for its remaining dependencies. Fetch them at execution. + 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]; + if (!(flags & (1U << 23))) continue; + const bool add = (flags & (1U << 26)) != 0; + if ((flags & (0x3bU << 24)) || (flags >> 30) != 1 || + (add && (flags & (7U << 20)))) return "unsupported QMD release format"; + if (add) { + if constexpr (!std::is_invocable_r_v<bool, Release, Address, Word, Word>) + return "unsupported QMD reduction callback"; + } + } + 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]; + if (!(flags & (1U << 23))) continue; + const Address destination = (Address{flags & 0xff} << 32) | qmd[24 + slot * 4]; + bool released{}; + if constexpr (std::is_invocable_r_v<bool, Release, Address, Word, Word>) { + released = release(destination, qmd[26 + slot * 4], flags); + } else { + released = release(destination, qmd[26 + slot * 4]); + } + if (!released) return "unmapped QMD release"; + } + if (qmd[16] & 1) { + const Word action = (qmd[16] >> 1) & 7; + if (action != 4 && action != 3 && action != 1) return "unsupported dependent QMD action"; + if (count == pending.size()) return "QMD chain limit"; + pending[count++] = {Address{qmd[15]} << 8, action == 4 ? 7U : action == 3 ? 3U : 2U, 1}; + } + if (qmd[16] & (1U << 5)) return "second dependent QMD unsupported"; + } + return {}; + } +private: + struct State { + Descriptor descriptor{}; + std::int32_t dependencies{}; + std::int32_t deferred{}; + bool loaded{}; + bool scheduled{}; + bool completed{}; + }; + std::unordered_map<Address, State> states; +}; +} // namespace Tegra diff --git a/src/video_core/native_query_trace.h b/src/video_core/native_query_trace.h new file mode 100644 index 0000000000..eae258dc04 --- /dev/null +++ b/src/video_core/native_query_trace.h @@ -0,0 +1,42 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include <cstdlib> +#include <fstream> +#include <iomanip> +#include <mutex> +#include "common/common_types.h" + +namespace VideoCommon { +namespace NativeQueryDiagnostics { +struct State { + State() : file{"native-query-trace.csv", std::ios::trunc} { + file << "event,asid,gpu,cpu,type,value,flags\n"; + } + std::mutex mutex; + std::ofstream file; + size_t records{}; +}; +inline State& GetState() { + static State state; + return state; +} +} // namespace NativeQueryDiagnostics + +// Events: 0=request, 1=written, 2=invalidated, 3=unsynchronized, +// 4=channel suspended, 5=channel resumed, 6=buffer readback, +// 7/8=QMD release slots observed, 9=PCAS request, +// 10=DMA destination (value=source, flags=line bytes), 11=inline write. +// Keep this opt-in and bounded. +inline void TraceNativeQuery(u32 event, u64 asid, u64 gpu, u64 cpu, u32 type, u64 value, u32 flags) { + if (!std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) return; + auto& state = NativeQueryDiagnostics::GetState(); + const std::scoped_lock lock{state.mutex}; + if (state.records++ < 1000000 || event == 4 || event == 5) { + state.file << event << ',' << asid << ',' << std::hex << gpu << ',' << cpu << ',' + << type << ',' << value << ',' << flags << std::dec << '\n'; + } + if (event == 4 || event == 5) state.file.flush(); +} +} // namespace VideoCommon diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.cpp b/src/video_core/renderer_vulkan/vk_texture_cache.cpp index 583cb6edaf..29afaac01b 100644 --- a/src/video_core/renderer_vulkan/vk_texture_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_texture_cache.cpp @@ -2085,6 +2085,11 @@ void Image::DownloadMemory(VkBuffer buffer, size_t offset, void Image::DownloadMemory(std::span<VkBuffer> buffers_span, std::span<size_t> offsets_span, std::span<const VideoCommon::BufferImageCopy> copies) { + if (!original_image) { + LOG_CRITICAL(Render_Vulkan, "Missing image at download: gpu={:#x}, cpu={:#x}, type={}, format={}, flags={:#x}, size={}, dimensions={}x{}", + gpu_addr, cpu_addr, static_cast<u32>(info.type), static_cast<u32>(info.format), + static_cast<u32>(flags), unswizzled_size_bytes, info.size.width, info.size.height); + } const bool is_rescaled = True(flags & ImageFlagBits::Rescaled); if (is_rescaled) { ScaleDown(); 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<u64, 6>& unique_hashes) { +bool ShaderCache::RefreshStages(std::array<u64, 6>& 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<u64, 6>& 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<u64, 6>& unique_hashes) { unique_hashes[index] = 0; continue; } - const GPUVAddr shader_addr{base_addr + shader_config.offset}; + const GPUVAddr shader_addr{native + ? (static_cast<u64>(shader_config.reserved[0] & 0xff) << 32) | shader_config.reserved[1] + : base_addr + shader_config.offset}; const std::optional<VAddr> 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<u64, 6>& 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<u64, 6>& 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<u64, NUM_PROGRAMS>& 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<Tegra::Engines::Maxwell3D::Regs::ShaderType>(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<u64>(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; diff --git a/src/video_core/shader_environment.cpp b/src/video_core/shader_environment.cpp index 4d5db7e511..fc2a26a568 100644 --- a/src/video_core/shader_environment.cpp +++ b/src/video_core/shader_environment.cpp @@ -5,7 +5,11 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include <algorithm> +#include <chrono> +#include <cstdlib> +#include <thread> #include <filesystem> +#include <cstring> #include <fstream> #include <memory> #include <optional> @@ -21,6 +25,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 +158,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 +188,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 +321,125 @@ 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) + throw Shader::RuntimeError("Native shader branch outside capture bounds"); + if (status == Shader::Ampere::ProgramCaptureState::Result::Complete) break; + } + auto lowered = Shader::Ampere::LowerToMaxwell(native); + // 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 +447,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 +497,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 +582,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 +682,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 +697,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); diff --git a/src/video_core/shader_environment.h b/src/video_core/shader_environment.h index 924e61ef46..e49f01d95f 100644 --- a/src/video_core/shader_environment.h +++ b/src/video_core/shader_environment.h @@ -47,6 +47,8 @@ public: [[nodiscard]] u64 ReadInstruction(u32 address) final; + [[nodiscard]] u32 TranslateProgramAddress(u32 address) const override; + [[nodiscard]] std::optional<u64> Analyze(); void SetCachedSize(size_t size_bytes); @@ -57,6 +59,10 @@ public: [[nodiscard]] size_t ReadSizeBytes() const noexcept; + [[nodiscard]] size_t SourceSizeBytes() const noexcept { + return native_source_size; + } + [[nodiscard]] bool CanBeSerialized() const noexcept; [[nodiscard]] u64 CalculateHash() const; @@ -72,6 +78,9 @@ public: protected: std::optional<u64> TryFindSize(); + void InitializeNativeCode(u32 source_header_size = 128); + size_t native_source_size{}; + Tegra::Texture::TICEntry ReadTextureInfo(GPUVAddr tic_addr, u32 tic_limit, bool via_header_index, u32 raw); @@ -172,6 +181,8 @@ public: [[nodiscard]] u64 ReadInstruction(u32 address) override; + [[nodiscard]] u32 TranslateProgramAddress(u32 address) const override; + [[nodiscard]] u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override; [[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override; diff --git a/src/video_core/texture_cache/accelerated_swizzle.cpp b/src/video_core/texture_cache/accelerated_swizzle.cpp index 4c3f724d79..8ae0ed218d 100644 --- a/src/video_core/texture_cache/accelerated_swizzle.cpp +++ b/src/video_core/texture_cache/accelerated_swizzle.cpp @@ -37,6 +37,7 @@ BlockLinearSwizzle2DParams MakeBlockLinearSwizzle2DParams(const SwizzleParameter .x_shift = GOB_SIZE_SHIFT + block.height + block.depth, .block_height = block.height, .block_height_mask = (1U << block.height) - 1, + .native_tiling = info.native_tiling, }; } @@ -63,6 +64,7 @@ BlockLinearSwizzle3DParams MakeBlockLinearSwizzle3DParams(const SwizzleParameter .block_height_mask = (1U << block.height) - 1, .block_depth = block.depth, .block_depth_mask = (1U << block.depth) - 1, + .native_tiling = info.native_tiling, }; } diff --git a/src/video_core/texture_cache/accelerated_swizzle.h b/src/video_core/texture_cache/accelerated_swizzle.h index d4250da689..8296950c38 100644 --- a/src/video_core/texture_cache/accelerated_swizzle.h +++ b/src/video_core/texture_cache/accelerated_swizzle.h @@ -20,6 +20,7 @@ struct BlockLinearSwizzle2DParams { u32 x_shift; u32 block_height; u32 block_height_mask; + u32 native_tiling; }; struct BlockLinearSwizzle3DParams { @@ -33,6 +34,7 @@ struct BlockLinearSwizzle3DParams { u32 block_height_mask; u32 block_depth; u32 block_depth_mask; + u32 native_tiling; }; [[nodiscard]] BlockLinearSwizzle2DParams MakeBlockLinearSwizzle2DParams( diff --git a/src/video_core/texture_cache/format_lookup_table.cpp b/src/video_core/texture_cache/format_lookup_table.cpp index d6e019cddf..20bb96aae3 100644 --- a/src/video_core/texture_cache/format_lookup_table.cpp +++ b/src/video_core/texture_cache/format_lookup_table.cpp @@ -157,6 +157,7 @@ PixelFormat PixelFormatFromTextureInfo(TextureFormat format, ComponentType red, return PixelFormat::S8_UINT_D24_UNORM; case Hash(TextureFormat::G24R8, UINT, UNORM, UNORM, UNORM, LINEAR): return PixelFormat::S8_UINT_D24_UNORM; + case Hash(TextureFormat::S8Z24, UNORM, UINT, UNORM, UNORM, LINEAR): case Hash(TextureFormat::S8Z24, UNORM, UINT, UINT, UINT, LINEAR): return PixelFormat::D24_UNORM_S8_UINT; case Hash(TextureFormat::Z32_X24S8, FLOAT, UINT, UNORM, UNORM, LINEAR): diff --git a/src/video_core/texture_cache/image_info.h b/src/video_core/texture_cache/image_info.h index eb490a6424..7407062d19 100644 --- a/src/video_core/texture_cache/image_info.h +++ b/src/video_core/texture_cache/image_info.h @@ -14,6 +14,17 @@ namespace VideoCommon { using Tegra::Texture::TICEntry; using VideoCore::Surface::PixelFormat; +// Native descriptors address the selected view directly. The legacy NVN +// layer-base bookkeeping is not an additional array prefix for this ABI. +inline TICEntry TextureViewHeader(TICEntry entry, bool native) { + if (native) { + entry.layer_base_0_2.Assign(0); + entry.layer_base_3_7.Assign(0); + entry.layer_base_8_10.Assign(0); + } + return entry; +} + struct ImageInfo { ImageInfo() = default; explicit ImageInfo(const TICEntry& config) noexcept; @@ -37,6 +48,7 @@ struct ImageInfo { u32 maybe_unaligned_layer_stride = 0; u32 num_samples = 1; u32 tile_width_spacing = 0; + bool native_tiling = false; bool rescaleable = false; bool downscaleable = false; bool forced_flushed = false; diff --git a/src/video_core/texture_cache/pending_download.h b/src/video_core/texture_cache/pending_download.h new file mode 100644 index 0000000000..52a48e82d3 --- /dev/null +++ b/src/video_core/texture_cache/pending_download.h @@ -0,0 +1,58 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later +#pragma once + +#include <algorithm> +#include <memory> +#include <unordered_set> +#include <cstddef> +#include <utility> +#include "common/slot_vector.h" + +namespace VideoCommon { + +template <class Image> +struct PendingImageDownload { + bool is_swizzle; + size_t async_buffer_id; + Common::SlotId object_id; + std::shared_ptr<Image> retired_image{}; + + Image& GetImage(Common::SlotVector<Image>& images) const { + return retired_image ? *retired_image : images[object_id]; + } +}; + +// Whole-image copies are recorded together at commit time. Repeated requests +// for the same live image therefore observe the same contents. Keep the first +// occurrence so reverse CPU writeback order remains unchanged for aliases. +template <class Image, class Downloads> +size_t DeduplicatePendingImageDownloads(Downloads& downloads, Common::SlotVector<Image>& images) { + std::unordered_set<const Image*> seen; + const size_t previous_size = downloads.size(); + const auto end = std::remove_if(downloads.begin(), downloads.end(), [&](const auto& download) { + return download.is_swizzle && !seen.insert(&download.GetImage(images)).second; + }); + downloads.erase(end, downloads.end()); + return previous_size - downloads.size(); +} + +// Cache slots may be recycled before queued GPU downloads are submitted or +// consumed. Keep the actual image, including its metadata, until those finish. +template <class Image, class Downloads, class Batches> +bool RetainPendingImageDownloads(Common::SlotId id, Image& image, Downloads& pending, + Batches& submitted) { + std::shared_ptr<Image> retained; + const auto retain = [&](auto& downloads) { + for (auto& download : downloads) { + if (!download.is_swizzle || download.object_id != id || download.retired_image) continue; + if (!retained) retained = std::make_shared<Image>(std::move(image)); + download.retired_image = retained; + } + }; + retain(pending); + for (auto& downloads : submitted) retain(downloads); + return retained != nullptr; +} + +} // namespace VideoCommon diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index e92e54f8ce..cd4ef82dbe 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -803,6 +803,17 @@ std::pair<typename P::ImageView*, bool> TextureCache<P>::TryFindFramebufferImage if (it == page_table.end()) { return {}; } + const auto view_format = [&]() { + switch (config.pixel_format) { + case Service::android::PixelFormat::Rgb565: + return PixelFormat::R5G6B5_UNORM; + case Service::android::PixelFormat::Bgra8888: + return PixelFormat::B8G8R8A8_UNORM; + default: + return PixelFormat::A8B8G8R8_UNORM; + } + }(); + const auto& image_map_ids = it->second; boost::container::small_vector<ImageId, 4> valid_image_ids; for (const ImageMapId map_id : image_map_ids) { @@ -814,20 +825,17 @@ std::pair<typename P::ImageView*, bool> TextureCache<P>::TryFindFramebufferImage if (image.image_view_ids.empty()) { continue; } + if (!VideoCore::Surface::IsViewCompatible(image.info.format, view_format, false, true)) { + static u32 rejected_alias_count{}; + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && rejected_alias_count++ < 8) { + LOG_WARNING(HW_GPU, "Ignoring incompatible framebuffer alias: address={:#x}, image={}, view={}", + cpu_addr, image.info.format, view_format); + } + continue; + } valid_image_ids.push_back(map.image_id); } - const auto view_format = [&]() { - switch (config.pixel_format) { - case Service::android::PixelFormat::Rgb565: - return PixelFormat::R5G6B5_UNORM; - case Service::android::PixelFormat::Bgra8888: - return PixelFormat::B8G8R8A8_UNORM; - default: - return PixelFormat::A8B8G8R8_UNORM; - } - }(); - const auto GetImageViewForFramebuffer = [&](ImageId image_id) { ImageViewInfo info{ImageViewType::e2D, view_format}; if (config.blending == Tegra::BlendMode::Opaque) { @@ -869,6 +877,7 @@ void TextureCache<P>::CommitAsyncFlushes() { // This is intentionally passing the value by copy if constexpr (IMPLEMENTS_ASYNC_DOWNLOADS) { auto& download_ids = uncommitted_downloads; + const size_t redundant_images = DeduplicatePendingImageDownloads(download_ids, slot_images); if (download_ids.empty()) { committed_downloads.emplace_back(std::move(uncommitted_downloads)); uncommitted_downloads.clear(); @@ -882,17 +891,32 @@ void TextureCache<P>::CommitAsyncFlushes() { for (PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { total_size_bytes += - Common::AlignUp(slot_images[download_info.object_id].unswizzled_size_bytes, 64); + Common::AlignUp(download_info.GetImage(slot_images).unswizzled_size_bytes, 64); any_none_dma = true; download_info.async_buffer_id = last_async_buffer_id; } } if (any_none_dma) { + if (std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS") && total_size_bytes >= 64_MiB) { + LOG_WARNING(HW_GPU, "Large image readback batch: bytes={}, entries={}, duplicates_removed={}, submitted_batches={}, retained_buffers={}", + total_size_bytes, download_ids.size(), redundant_images, + committed_downloads.size(), async_buffers_death_ring.size()); + size_t traced_images = 0; + for (const auto& download : download_ids) { + if (!download.is_swizzle || traced_images++ >= 8) continue; + const auto& image = download.GetImage(slot_images); + LOG_WARNING(HW_GPU, "Readback image: gpu={:#x}, cpu={:#x}, format={}, size={}x{}x{}, levels={}, layers={}, layer_stride={}, guest_bytes={}, download_bytes={}", + image.gpu_addr, image.cpu_addr, image.info.format, + image.info.size.width, image.info.size.height, image.info.size.depth, + image.info.resources.levels, image.info.resources.layers, + image.info.layer_stride, image.guest_size_bytes, image.unswizzled_size_bytes); + } + } auto download_map = runtime.DownloadStagingBuffer(total_size_bytes, true); for (const PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { - Image& image = slot_images[download_info.object_id]; + Image& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); image.DownloadMemory(download_map, copies); download_map.offset += Common::AlignUp(image.unswizzled_size_bytes, 64); @@ -925,7 +949,7 @@ void TextureCache<P>::PopAsyncFlushes() { auto& download_info = download_ids[i - 1]; auto& download_buffer = download_map[download_info.async_buffer_id]; if (download_info.is_swizzle) { - const ImageBase& image = slot_images[download_info.object_id]; + const ImageBase& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); download_buffer.offset -= Common::AlignUp(image.unswizzled_size_bytes, 64); std::span<u8> download_span = @@ -954,7 +978,7 @@ void TextureCache<P>::PopAsyncFlushes() { size_t total_size_bytes = 0; for (const PendingDownload& download_info : download_ids) { if (download_info.is_swizzle) { - total_size_bytes += slot_images[download_info.object_id].unswizzled_size_bytes; + total_size_bytes += download_info.GetImage(slot_images).unswizzled_size_bytes; } } auto download_map = runtime.DownloadStagingBuffer(total_size_bytes); @@ -963,7 +987,7 @@ void TextureCache<P>::PopAsyncFlushes() { if (!download_info.is_swizzle) { continue; } - Image& image = slot_images[download_info.object_id]; + Image& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); image.DownloadMemory(download_map, copies); download_map.offset += image.unswizzled_size_bytes; @@ -976,7 +1000,7 @@ void TextureCache<P>::PopAsyncFlushes() { if (!download_info.is_swizzle) { continue; } - const ImageBase& image = slot_images[download_info.object_id]; + const ImageBase& image = download_info.GetImage(slot_images); const auto copies = FullDownloadCopies(image.info); SwizzleImage(*gpu_memory, image.gpu_addr, image.info, copies, download_span, swizzle_data_buffer); download_map.offset += image.unswizzled_size_bytes; @@ -1178,7 +1202,18 @@ void TextureCache<P>::UploadImageContents(Image& image, StagingBuffer& staging) } template <class P> -ImageViewId TextureCache<P>::CreateImageView(const TICEntry& config) { +ImageViewId TextureCache<P>::CreateImageView(const TICEntry& raw_config) { + const bool native = gpu_memory->UsesNativeTiling(); + if (native && raw_config.BaseLayer() && std::getenv("EDEN_NATIVE_DRAW_DIAGNOSTICS")) { + static u32 texture_view_traces{}; + if (texture_view_traces++ < 8) { + LOG_WARNING(HW_GPU, "Native texture view header: address={:#x}, dimensions={}x{}x{}, type={}, legacy_layer={}, raw={:016x}/{:016x}/{:016x}/{:016x}", + raw_config.Address(), raw_config.Width(), raw_config.Height(), raw_config.Depth(), + static_cast<u32>(raw_config.texture_type.Value()), raw_config.BaseLayer(), + raw_config.raw[0], raw_config.raw[1], raw_config.raw[2], raw_config.raw[3]); + } + } + const TICEntry config = TextureViewHeader(raw_config, native); const ImageInfo info(config); if (info.type == ImageType::Buffer) { const ImageViewInfo view_info(config, 0); @@ -1526,7 +1561,9 @@ ImageId TextureCache<P>::InsertImage(const ImageInfo& info, GPUVAddr gpu_addr, } } ASSERT_MSG(cpu_addr, "Tried to insert an image to an invalid gpu_addr={:#x}", gpu_addr); - const ImageId image_id = JoinImages(info, gpu_addr, *cpu_addr); + ImageInfo storage_info = info; + storage_info.native_tiling = gpu_memory->UsesNativeTiling(); + const ImageId image_id = JoinImages(storage_info, gpu_addr, *cpu_addr); const Image& image = slot_images[image_id]; // Using "image.gpu_addr" instead of "gpu_addr" is important because it might be different const auto [it, is_new] = image_allocs_table.try_emplace(image.gpu_addr); @@ -2385,7 +2422,11 @@ void TextureCache<P>::DeleteImage(ImageId image_id, bool immediate_delete) { } slot_image_views.erase(image_view_id); } - if (!immediate_delete) { + const bool retained = RetainPendingImageDownloads(image_id, slot_images[image_id], + uncommitted_downloads, committed_downloads); + if (retained) { + LOG_INFO(Render_Vulkan, "Retained image for queued download: id={}", image_id.Value()); + } else if (!immediate_delete) { sentenced_images.Push(std::move(slot_images[image_id])); } slot_images.erase(image_id); diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h index 8a2318cc72..d103d66189 100644 --- a/src/video_core/texture_cache/texture_cache_base.h +++ b/src/video_core/texture_cache/texture_cache_base.h @@ -28,6 +28,7 @@ #include <ranges> #include "common/scratch_buffer.h" #include "common/slot_vector.h" +#include "video_core/texture_cache/pending_download.h" #include "common/thread_worker.h" #include "video_core/compatible_formats.h" #include "video_core/control/channel_state_cache.h" @@ -463,11 +464,7 @@ private: size_t size; }; - struct PendingDownload { - bool is_swizzle; - size_t async_buffer_id; - Common::SlotId object_id; - }; + using PendingDownload = PendingImageDownload<Image>; Common::SlotVector<Image> slot_images; Common::SlotVector<ImageMapView> slot_map_views; diff --git a/src/video_core/texture_cache/util.cpp b/src/video_core/texture_cache/util.cpp index e55d0752ec..21070b2d25 100644 --- a/src/video_core/texture_cache/util.cpp +++ b/src/video_core/texture_cache/util.cpp @@ -439,7 +439,7 @@ template <u32 GOB_EXTENT> } layers = 1; } else { - layers = (std::max)(resources.layers, info.resources.layers + base->layer); + layers = (std::max)(info.resources.layers, resources.layers + base->layer); } return OverlapResult{ .gpu_addr = overlap.gpu_addr, @@ -566,7 +566,7 @@ void SwizzleBlockLinearImage(Tegra::MemoryManager& gpu_memory, GPUVAddr gpu_addr dst(gpu_memory, gpu_addr + guest_offset, subresource_size, &tmp_buffer); SwizzleTexture(dst, src, bytes_per_block, num_tiles.width, num_tiles.height, - num_tiles.depth, block.height, block.depth); + num_tiles.depth, block.height, block.depth, 1, gpu_memory.UsesNativeTiling()); } host_offset += host_bytes_per_layer; @@ -913,7 +913,7 @@ boost::container::small_vector<BufferImageCopy, 16> UnswizzleImage(Tegra::Memory const std::span<u8> dst = output.subspan(host_offset); const std::span<const u8> src = input.subspan(guest_offset + guest_layer_offset); UnswizzleTexture(dst, src, 1U << bpp_log2, num_tiles.width, num_tiles.height, - num_tiles.depth, block.height, block.depth, stride_alignment); + num_tiles.depth, block.height, block.depth, stride_alignment, gpu_memory.UsesNativeTiling()); guest_layer_offset += layer_stride; host_offset += host_bytes_per_layer; } |