From 41e7291f794ce5049c1de6f1b1f08db9b07d215c Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:19:22 +0100 Subject: [PATCH 01/11] shader: lower Ocarina ISETP.EX accumulation and F32 global reductions Ported from zortos293/Lichii 3aa89d1cdb (fix/oot-bloom). ISETP.EX now computes into a free temporary predicate and merges it with the accumulate predicate and second destination, instead of rejecting those forms. RED takes F32 ADD and a signed offset on the register address. High-register ALD/FADD/FMUL/FFMA operands get shadow copies. Not ported: the routed SPIR-V GlobalAtomicUMax32 (Eglantine has no routed global atomic functions yet) and the native compute SSBO sizing (already present from 2e12e0def5). Assisted-by: Claude Opus 5.5 (ported the upstream patch hunks and resolved conflicts) (cherry picked from commit 350b7d8ffd85eccc638ffa8751ce31c3ba502c57) --- .../frontend/ampere/lower_to_maxwell.h | 58 ++++++++++++------- 1 file changed, 37 insertions(+), 21 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 66748f2bcd..c87d36cd98 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -81,6 +81,9 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re case 0x308: operand(16); operand(32); break; case 0x433: case 0x633: operand(16); operand(24); break; case 0x233: operand(16); operand(24); operand(32); break; + case 0xa20: case 0x621: operand(16); operand(24); break; + case 0xa23: operand(16); operand(24); operand(64); break; + case 0x321: operand(16, input.Bits(74, 2) + 1); operand(24); operand(32); break; case 0x812: case 0x824: case 0x831: case 0x424: case 0x624: operand(16); operand(24); operand(64); break; case 0x230: case 0x231: case 0x211: case 0x224: @@ -1138,21 +1141,22 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 break; } case 0x98e: { // RED.E.{ADD..XOR}.STRONG.GPU U32 (op in bits 87-89), register address. - if ((insn.words[2] & ~0x0380003fU) == 0x0c10e100 && insn.Bits(24, 8) == 255) { - // [URa + offset]: a uniform 64-bit address with a signed byte offset. - const u32 address = uniform_register(insn.Bits(64, 6)); - const s32 offset = static_cast(insn.Bits(40, 24) << 8) >> 8; - if (address == ~0U || (address & 1) || offset < -0x80000 || offset > 0x7ffff) - return fail(index, "native global reduction uniform address"); - emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | (u64{address} << 8) | - (u64{insn.Bits(87, 3)} << 23) | (u64{static_cast(offset) & 0xfffff} << 28) | - (u64{1} << 48)); - break; - } - if ((insn.words[2] & ~0x03800000U) != 0x0010e100 || insn.Bits(40, 24) || (insn.Bits(24, 8) & 1)) + const u32 type = insn.Bits(73, 3); + const u32 operation = insn.Bits(87, 3); + const bool uniform_address = insn.Bits(24, 8) == 255; + const u32 variable = 0x03800e00U | (uniform_address ? 0x3fU : 0U); + const u32 mode = uniform_address ? 0x0c10e100U : 0x0010e100U; + const s32 offset = static_cast(insn.Bits(40, 24) << 8) >> 8; + if ((insn.words[2] & ~variable) != mode || + (type != 0 && type != 3) || (type == 3 && operation != 0)) return fail(index, "native global reduction mode"); - emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | - (u64{insn.Bits(24, 8)} << 8) | (u64{insn.Bits(87, 3)} << 23) | (u64{1} << 48)); + const u32 address = uniform_address ? uniform_register(insn.Bits(64, 6)) + : insn.Bits(24, 8); + if (address == ~0U || (address & 1) || offset < -0x80000 || offset > 0x7ffff) + return fail(index, "native global reduction address"); + emit(0xebf8000000000000ULL | predicate | insn.Bits(32, 8) | (u64{address} << 8) | + (u64{type} << 20) | (u64{operation} << 23) | + (u64{static_cast(offset) & 0xfffff} << 28) | (u64{1} << 48)); break; } case 0x2bd: // UFLO.U32, the highest set-bit index (.SH, bit 74: its shift amount). @@ -2709,18 +2713,25 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 // result comes in as the predicate at bits 68-70 (negated by 71), e.g. // ISETP.NE P0 = lo != 0, then ISETP.NE.EX P0 = (hi != 0) OR P0. EQ: high equal // AND low; NE: high different OR low; the orderings: high strictly ordered OR - // (high equal AND low). One or two Maxwell ISETPs into the destination; only - // the plain form (no second destination, AND with PT) is handled. + // (high equal AND low). const u32 compare = insn.Bits(76, 3); const u32 destination = insn.Bits(81, 3); + const u32 destination_b = insn.Bits(84, 3); const u32 low = insn.Bits(68, 4); - if (insn.Bits(84, 3) != 7 || insn.Bits(87, 4) != 7 || insn.Bits(74, 2) != 0 || - compare == 0 || compare == 7 || - (destination != 7 && insn.Bits(12, 3) == destination)) + const u32 accumulate = insn.Bits(87, 4); + const u32 combine = insn.Bits(74, 2); + if (compare == 0 || compare == 7 || combine == 3) return fail(index, "native extended integer comparison mode"); const u32 a = operand(0, false), b = operand(1, false); + u32 temporary = 0; + while (temporary == destination || temporary == destination_b || + temporary == (low & 7) || temporary == (accumulate & 7) || + temporary == insn.Bits(12, 3)) + ++temporary; + const u32 mask = 1U << temporary; + emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); const auto isetp = [&](u32 op, u32 input, u32 combine) { - return 0x5b60000000000000ULL | predicate | 7 | (u64{destination} << 3) | + return 0x5b60000000000000ULL | predicate | 7 | (u64{temporary} << 3) | (u64{a} << 8) | (u64{b} << 20) | (u64{input & 7} << 39) | (u64{input >> 3} << 42) | (u64{combine} << 45) | (u64{insn.Bits(73, 1)} << 48) | (u64{op} << 49); @@ -2732,8 +2743,13 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 emit(isetp(Ne, low, Or)); } else { emit(isetp(Eq, low, And)); - emit(isetp(compare <= 3 ? Lt : Gt, destination, Or)); + emit(isetp(compare <= 3 ? Lt : Gt, temporary, Or)); } + emit(0x5090000000000000ULL | predicate | destination_b | + (u64{destination} << 3) | (u64{temporary} << 12) | (u64{7} << 29) | + (u64{accumulate & 7} << 39) | (u64{accumulate >> 3} << 42) | + (u64{combine} << 45)); + emit(0x38f0000000070000ULL | (u64{253} << 8) | (u64{mask} << 20)); break; } const u32 a = operand(0, floating), b = operand(1, floating); From 245d902e98649b452650d04b906b213aae225190 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:22:15 +0100 Subject: [PATCH 02/11] shader: lower more Ocarina compute and high-register graphics forms Ported from zortos293/Lichii 35c043b3ef (fix/oot-bloom). Adds uniform IMAD with an immediate addend (UIMAD 0x4a4), signed REDUX.MIN/MAX, S32 RED, negated scalar F2I, packed F2IP.U8.F32 with an RZ second source, SHF.R.U64 by an immediate, and shadow operands for the remaining bloom-capture opcodes. Eglantine's F2I has no unsigned byte form, so F2I.U8 (which F2IP lowers to) gets the float clamp to [0, 255] that Lichii's F2I uses. The bounded-control write_count table is not ported; Eglantine does not have that call proof. Assisted-by: Claude Opus 5.5 (ported the upstream patch hunks and resolved conflicts) (cherry picked from commit 3f065ef733e3c2b889ccd3ad54705ec6af8b13dd) --- .../frontend/ampere/lower_to_maxwell.h | 72 +++++++++++++++---- 1 file changed, 59 insertions(+), 13 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index c87d36cd98..004cd7f9cb 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -76,14 +76,18 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re const u32 opcode = input.Bits(0, 12); switch (opcode) { case 0x802: case 0xa02: case 0xb08: case 0xd08: operand(16); break; - case 0x804: operand(24); break; + case 0x804: case 0x80c: operand(24); break; case 0x805: operand(16, 2); break; case 0x308: operand(16); operand(32); break; case 0x433: case 0x633: operand(16); operand(24); break; case 0x233: operand(16); operand(24); operand(32); break; - case 0xa20: case 0x621: operand(16); operand(24); break; - case 0xa23: operand(16); operand(24); operand(64); break; + case 0xa20: case 0x621: case 0x808: case 0x923: operand(16); operand(24); break; + case 0x209: operand(16); operand(24); operand(32); break; + case 0xa23: case 0x623: operand(16); operand(24); operand(64); break; case 0x321: operand(16, input.Bits(74, 2) + 1); operand(24); operand(32); break; + case 0x322: operand(24); operand(32, input.Bits(74, 2) + 1); break; + case 0x927: operand(24); operand(32); break; + case 0x202: case 0x243: operand(16); operand(32); break; case 0x812: case 0x824: case 0x831: case 0x424: case 0x624: operand(16); operand(24); operand(64); break; case 0x230: case 0x231: case 0x211: case 0x224: @@ -183,6 +187,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 }); for (const auto& instruction : code) { const u32 op = instruction.Bits(0, 12); + uses_uniform_bank |= op == 0x4a4; uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0x896 || op == 0xf60 || op == 0xc11 || op == 0xd08 || 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 == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || op == 0x8a4 || op == 0xd05 || (op == 0x984 && instruction.Bits(91, 1)) || (op == 0x988 && instruction.Bits(91, 1)) || op == 0x89c || op == 0x883 || op == 0x2a5 || op == 0xd06 || op == 0xf63 || op == 0xf66 || op == 0xf69 || op == 0xc0b || op == 0x8a5 || op == 0x299 || op == 0xc09 || (op == 0x81c && instruction.Bits(67, 1)) || op == 0xc0a || (op == 0x83b && instruction.Bits(91, 1)); @@ -280,7 +285,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 bool rewritten_uniform = false; bool uniform_guard = false; u32 uniform_output = 7; // ULOP3's UPn output: computed in P(6-n), then merged into R239 - if (opcode == 0x2a4 || opcode == 0x8a4 || opcode == 0x891 || opcode == 0x890 || opcode == 0x290 || + if (opcode == 0x2a4 || opcode == 0x4a4 || opcode == 0x8a4 || opcode == 0x891 || opcode == 0x890 || opcode == 0x290 || opcode == 0x292 || opcode == 0x892 || opcode == 0x291 || opcode == 0x896 || opcode == 0x8a5) { // A guarded ULEA without .HI or a carry out lowers to one guarded ISCADD, which // leaves R253 to hold the guest predicate (Pokemon Pokopia's fullscreen passes). @@ -316,7 +321,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 uniform_output = insn.Bits(81, 3); insn.words[2] = (insn.words[2] & ~(7U << 17)) | ((6 - uniform_output) << 17); } - opcode = opcode == 0x2a4 ? 0x224 : opcode == 0x8a4 ? 0x824 : opcode == 0x891 ? 0x811 : opcode == 0x890 ? 0x810 + opcode = opcode == 0x2a4 ? 0x224 : opcode == 0x4a4 ? 0x424 : opcode == 0x8a4 ? 0x824 : opcode == 0x891 ? 0x811 : opcode == 0x890 ? 0x810 : opcode == 0x292 ? 0x212 : opcode == 0x892 ? 0x812 : opcode == 0x291 ? 0x211 : opcode == 0x896 ? 0x816 : opcode == 0x8a5 ? 0x825 : 0x210; } @@ -1148,7 +1153,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const u32 mode = uniform_address ? 0x0c10e100U : 0x0010e100U; const s32 offset = static_cast(insn.Bits(40, 24) << 8) >> 8; if ((insn.words[2] & ~variable) != mode || - (type != 0 && type != 3) || (type == 3 && operation != 0)) + (type != 0 && type != 1 && type != 3) || (type == 3 && operation != 0)) return fail(index, "native global reduction mode"); const u32 address = uniform_address ? uniform_register(insn.Bits(64, 6)) : insn.Bits(24, 8); @@ -1215,17 +1220,23 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 } case 0x3c4: { // REDUX.SUM (modulo 2^32), .MIN/.MAX (unsigned), .OR or .XOR over the actual active lanes. const u32 source = insn.Bits(24, 8), target = uniform_register(static_cast(dest)); - const bool maximum = insn.words[2] == 0x14000, minimum = insn.words[2] == 0x10000; + const u32 mode = insn.words[2] & ~0x200U; + const bool maximum = mode == 0x14000, minimum = mode == 0x10000; + const bool signed_order = insn.Bits(73, 1); // OR and XOR: inactive lanes contribute zero, as for the sum. Maxwell LOP: 1 = OR, 2 = XOR. const u32 logical = insn.words[2] == 0x4000 ? 1U : insn.words[2] == 0x8000 ? 2U : 0U; if (target == ~0U || target == 255 || insn.words[1] || - (insn.words[2] != 0xc000 && !maximum && !minimum && !logical) || native_predicate != 7) + (insn.words[2] != 0xc000 && !maximum && !minimum && !logical) || + (signed_order && !maximum && !minimum) || native_predicate != 7) return fail(index, "native warp sum mode"); emit(0x50d8000000070000ULL | 250 | (u64{7} << 39) | (u64{7} << 45) | (u64{1} << 48)); mov_imm(target, minimum ? ~0U : 0U); 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)); + if (signed_order) + emit(0x0400000000070000ULL | 251 | (u64{251} << 8) | + (u64{0x80000000U} << 20) | (u64{2} << 53)); 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)); @@ -1243,6 +1254,9 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 emit(0x5c10000000070000ULL | target | (u64{target} << 8) | (u64{251} << 20)); } } + if (signed_order) + emit(0x0400000000070000ULL | target | (u64{target} << 8) | + (u64{0x80000000U} << 20) | (u64{2} << 53)); break; } case 0xf8c: { // ATOMS.POPC.INC.32 with no returned value: one increment per active lane. @@ -1833,6 +1847,14 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(78, 2)} << 39)); break; } + case 0x243: // F2IP.U8.F32 with RZ for the second source: one unsigned byte, as F2I.U8 + if (insn.Bits(24, 8) != 255 || (insn.words[1] & ~0xffU) || + insn.words[2] != 0x0000c4ffU) + return fail(index, "native packed float-to-byte mode"); + insn.words[0] &= 0x00ffffffU; + insn.words[2] = 0x0021e000U; + opcode = 0x305; + [[fallthrough]]; case 0x305: case 0xb05: case 0xd05: { // F2I NTZ, register/constant-buffer/uniform source const u32 mode = insn.words[2] & ~(opcode == 0xd05 ? 0x08000000U : 0U); // F16 source: bit 72 signs the result as in the other forms, bits 60-61 pick H0/H1. @@ -1845,9 +1867,14 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const bool short_trunc = opcode == 0x305 && (mode & ~0x100U) == 0x0021e800; // Bit 62: |Rs| on the register form (Hyrule Warriors: F2I.FTZ.U16.NTZ R26, |R42|). const bool absolute = opcode == 0x305 && !half && insn.Bits(62, 1); + const bool negate = opcode == 0x305 && !half && insn.Bits(63, 1); + // F2I.U8: clamped to [0, 255] as a float, then converted as a 32-bit integer. + const bool unsigned_byte = opcode == 0x305 && (mode & ~0x10000U) == 0x0020e000; if ((!half && !small && !floor && !nearest && !round_up && !signed_byte && !short_trunc && - (mode & ~0x10100U) != 0x0020f000) || - (opcode != 0xb05 && (insn.words[1] & (half ? 0xefffff00U : absolute ? 0xbfffff00U : 0xffffff00U))) || + !unsigned_byte && (mode & ~0x10100U) != 0x0020f000) || + (unsigned_byte && (insn.Bits(24, 8) || (insn.words[3] & 0xc00001ffU))) || + (opcode != 0xb05 && (insn.words[1] & ~(half ? 0x100000ffU : + 0xffU | (absolute ? 0x40000000U : 0U) | (negate ? 0x80000000U : 0U)))) || (opcode == 0xd05 && uniform_register(insn.Bits(32, 8)) == ~0U)) { return fail(index, "native float-to-integer conversion mode"); } @@ -1858,8 +1885,20 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 extract_half({source, insn.Bits(60, 1) ? 3U : 2U, false, false}, 0, 252); source = 252; } - if (absolute) source = materialize(0, source, 249, true, false); - if (!is_signed) { + if (absolute || negate) source = materialize(0, source, 249, absolute, negate); + if (unsigned_byte) { + 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)); + mov_imm(252, 0); + emit(0x5c60000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{252} << 20) | (u64{15} << 39)); + mov_imm(252, 0x437f0000U); + emit(0x5c60000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{252} << 20) | (u64{7} << 39)); + source = 250; + } else 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) | @@ -1885,7 +1924,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 emit(base(0x5ce0000000000000ULL) | (u64{1} << 10) | (u64{1} << 12) | (u64{1} << 13) | (dest << 20) | (u64{1} << 50)); } - if (!is_signed && !small && !short_trunc) { + if (!is_signed && !small && !short_trunc && !unsigned_byte) { emit(base(0x5c40000000000000ULL) | (dest << 8) | (u64{251} << 20) | (u64{1} << 41) | (u64{7} << 48)); } @@ -2817,6 +2856,13 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 } 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(73, 2) == 1 && + !insn.Bits(75, 1) && insn.words[1] <= 32) { + emit(base(0x38f8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | + (u64{insn.words[1]} << 20) | (u64{2} << 37) | + (u64{insn.Bits(64, 8)} << 39)); + break; + } 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)); From 2d29137f253c6a51363033096e615b8b63f12598 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:23:43 +0100 Subject: [PATCH 03/11] shader: lower Ocarina CALL helpers, 64-bit shifts and high uniform forms Ported from zortos293/Lichii a4cc1e680f (fix/oot-bloom). Capture now treats CALL.REL.NOINC as a call, as the lowering already does, so a program whose helpers sit past the EXIT ends at the last RET instead of running into the next program. USHF and SHF.R.U64 lower to a full 64-bit right shift clamped at 64, and Maxwell SHR clamps an unsigned shift of 32 or more to zero instead of sign-filling. USEL takes a UPn guard. UBREV, UFLO and UPOPC map every uniform register, and UPOPC marks the program as using the uniform bank. The bounded-control call-proof changes are not ported; Eglantine does not have that analysis. Assisted-by: Claude Opus 5.5 (ported the upstream patch hunks and tests, resolved conflicts) (cherry picked from commit bb5998f4977d51ebd8506b87a1c4a4588237c85a) --- .../frontend/ampere/lower_to_maxwell.h | 77 +++++++++++++++---- .../frontend/ampere/program_capture.h | 9 +-- .../translate/impl/integer_shift_right.cpp | 5 +- src/tests/video_core/ampere_lowering.cpp | 70 +++++++++++++++++ 4 files changed, 137 insertions(+), 24 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 004cd7f9cb..fa90ec0642 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -83,14 +83,14 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re case 0x233: operand(16); operand(24); operand(32); break; case 0xa20: case 0x621: case 0x808: case 0x923: operand(16); operand(24); break; case 0x209: operand(16); operand(24); operand(32); break; - case 0xa23: case 0x623: operand(16); operand(24); operand(64); break; + case 0xa23: case 0x623: case 0x423: operand(16); operand(24); operand(64); break; case 0x321: operand(16, input.Bits(74, 2) + 1); operand(24); operand(32); break; case 0x322: operand(24); operand(32, input.Bits(74, 2) + 1); break; case 0x927: operand(24); operand(32); break; case 0x202: case 0x243: operand(16); operand(32); break; case 0x812: case 0x824: case 0x831: case 0x424: case 0x624: operand(16); operand(24); operand(64); break; - case 0x230: case 0x231: case 0x211: case 0x224: + case 0x230: case 0x231: case 0x211: case 0x212: case 0x224: operand(16); operand(24); operand(32); operand(64); break; case 0x819: case 0xa11: operand(16); operand(24); operand(64); break; @@ -187,7 +187,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 }); for (const auto& instruction : code) { const u32 op = instruction.Bits(0, 12); - uses_uniform_bank |= op == 0x4a4; + uses_uniform_bank |= op == 0x4a4 || op == 0x2bf; uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0x896 || op == 0xf60 || op == 0xc11 || op == 0xd08 || 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 == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || op == 0x8a4 || op == 0xd05 || (op == 0x984 && instruction.Bits(91, 1)) || (op == 0x988 && instruction.Bits(91, 1)) || op == 0x89c || op == 0x883 || op == 0x2a5 || op == 0xd06 || op == 0xf63 || op == 0xf66 || op == 0xf69 || op == 0xc0b || op == 0x8a5 || op == 0x299 || op == 0xc09 || (op == 0x81c && instruction.Bits(67, 1)) || op == 0xc0a || (op == 0x83b && instruction.Bits(91, 1)); @@ -483,6 +483,20 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 emit(base(0x5c18000000000000ULL) | (u64{251} << 8) | (u64{250} << 20) | (u64{16} << 39)); }; + const auto right_shift_u64_low = [&](u32 target, u32 low, u32 high, u32 shift) { + mov_imm(250, 32); + emit(0x5c10000000000000ULL | predicate | 251 | (u64{250} << 8) | + (u64{shift} << 20) | (u64{1} << 48)); + emit(0x5c28000000000000ULL | predicate | 250 | (u64{low} << 8) | (u64{shift} << 20)); + emit(0x5c48000000000000ULL | predicate | 251 | (u64{high} << 8) | (u64{251} << 20)); + emit(0x5c40000000000000ULL | predicate | 250 | (u64{250} << 8) | + (u64{251} << 20) | (u64{1} << 41) | (u64{7} << 48)); + emit(0x1c00000000000000ULL | predicate | 251 | (u64{shift} << 8) | + (u64{0xffffffe0U} << 20)); + emit(0x5c28000000000000ULL | predicate | 251 | (u64{high} << 8) | (u64{251} << 20)); + emit(0x5c40000000000000ULL | predicate | target | (u64{250} << 8) | + (u64{251} << 20) | (u64{1} << 41) | (u64{7} << 48)); + }; switch (opcode) { case 0x304: { // F2F.F16.F32, scalar round-to-nearest conversion. if (insn.words[2] != 0x00200800 || insn.Bits(24, 8) || @@ -1036,6 +1050,12 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const u32 high_reg = uniform_register(high); const bool dst_ok = dst != 63 && dst_reg != ~0U; const bool low_ok = low != 63 && low_reg != ~0U, high_ok = high != 63 && high_reg != ~0U; + if (dst_ok && low_reg != ~0U && high_reg != ~0U && + insn.words[2] == (0x08001200U | high)) { + mov_imm(252, insn.words[1]); + right_shift_u64_low(dst_reg, low_reg, high_reg, 252); + break; + } // A UPn guard (Hyrule Warriors: @UP1 USHF.R.U32.HI) is swapped into the predicate above. if (insn.words[2] == (0x08011400U | high) && low == 63 && high_ok && dst_ok && insn.words[1] <= 32) { @@ -1164,23 +1184,36 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{static_cast(offset) & 0xfffff} << 28) | (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) + case 0x2bd: { // UFLO.U32, the highest set-bit index (.SH, bit 74: its shift amount). + const u32 target = uniform_register(static_cast(dest)); + const u32 source = uniform_register(insn.Bits(32, 8)); + if (dest == 63 || target == ~0U || source == ~0U || (insn.words[1] & ~0xffU) || + (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) | + emit(0x5c30000000000000ULL | predicate | target | (u64{source} << 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) + } + case 0x2be: { // UBREV, as a bit-reversing zero shift like BREV. + const u32 target = uniform_register(static_cast(dest)); + const u32 source = uniform_register(insn.Bits(32, 8)); + if (dest == 63 || target == ~0U || source == ~0U || (insn.words[1] & ~0xffU) || + 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) | + emit(0x3828000000000000ULL | predicate | target | (u64{source} << 8) | (u64{1} << 39) | (u64{1} << 40)); break; - case 0x2bf: // UPOPC, as POPC of the uniform's register. - if (dest >= 16 || insn.words[1] >= 16 || insn.words[2] != 0x08000000 || native_predicate != 7) + } + case 0x2bf: { // UPOPC, as POPC of the uniform's register (bit 63 inverts it). + const u32 target = uniform_register(static_cast(dest)); + const u32 source = uniform_register(insn.Bits(32, 8)); + if (dest == 63 || target == ~0U || source == ~0U || (insn.words[1] & ~0x800000ffU) || + insn.words[2] != 0x08000000 || native_predicate != 7) return fail(index, "native uniform population count mode"); - emit(0x5c08000000000000ULL | predicate | (dest + 216) | (u64{insn.words[1] + 216} << 20)); + emit(0x5c08000000000000ULL | predicate | target | (u64{source} << 20) | + (u64{insn.Bits(63, 1)} << 40)); break; + } case 0x3a1: { // MATCH.ALL (.U32, or .U64 on bit 73) with only the all-equal predicate retained. const u32 source = insn.Bits(24, 8), target = insn.Bits(81, 3); const u32 words = insn.Bits(73, 1) ? 2 : 1; @@ -1373,6 +1406,11 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const bool shift_operand = constant ? valid_cbuf() && !insn.Bits(62, 2) : insn.words[1] <= 255; const auto load_shift = [&] { if (constant) materialize(2, 0, 251, false, false); }; const bool right = insn.Bits(76, 1) != 0, high_result = insn.Bits(80, 1) != 0; + if (opcode == 0x219 && shift_operand && + (insn.words[2] & ~0xffU) == 0x00001200U) { + right_shift_u64_low(static_cast(dest), low, high, shift); + break; + } // SHF.R.U32 of the low word with a zero high word: a plain logical right shift. if (right && !high_result && high == 255 && insn.Bits(73, 2) == 3 && shift != ~0U && shift_operand && @@ -1705,7 +1743,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const u32 target = uniform_register(static_cast(dest)); const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x887 ? 252U : uniform_register(insn.Bits(32, 8)); if (condition > 6 || target == ~0U || a == ~0U || b == ~0U || - insn.Bits(12, 4) != 7 || (insn.words[2] & ~0x07800000U) != 0x08000000U) + (insn.words[2] & ~0x07800000U) != 0x08000000U) return fail(index, "native uniform select mode"); // The clamp's UISETP may sit a few instructions back, with only instructions between // that write no uniform state and do not branch (Nintendo Switch Sports' jump tables), @@ -1723,7 +1761,8 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 } const bool target_between = std::any_of(relocations.begin(), relocations.end(), [&](const Relocation& relocation) { return relocation.target > compare && relocation.target <= index; }); - if (opcode == 0x887 && compare < index && !target_between && condition == 0 && insn.words[1] <= 0x7ffff && + if (opcode == 0x887 && native_predicate == 7 && compare < index && !target_between && + condition == 0 && insn.words[1] <= 0x7ffff && code[compare].Bits(12, 4) == 7 && code[compare].Bits(24, 8) == insn.Bits(24, 8) && code[compare].words[1] == insn.words[1] && code[compare].words[2] == 0x0bf01070 && insn.words[2] == 0x08000000) { @@ -1732,11 +1771,15 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (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; + if (opcode == 0x887) + emit(0x0100000000000000ULL | (u64{7} << 16) | 252 | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); + const u32 temporary = 6 - condition; + const u32 mask = (1U << temporary) | (native_predicate != 7 ? 1U << guard_slot : 0U); 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) | + emit(0x5ca0000000000000ULL | (u64{(native_predicate & 8) | (native_predicate != 7 ? guard_slot : 7U)} << 16) | + 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; diff --git a/src/shader_recompiler/frontend/ampere/program_capture.h b/src/shader_recompiler/frontend/ampere/program_capture.h index fe9653e2d7..985aac1e66 100644 --- a/src/shader_recompiler/frontend/ampere/program_capture.h +++ b/src/shader_recompiler/frontend/ampere/program_capture.h @@ -21,10 +21,9 @@ class ProgramCaptureState { const s64 target = s64{index} * 16 + 16 + relative; if (target < 0 || target % 16 || target >= MaximumInstructions * 16) return Result::InvalidTarget; - // CALL.REL.NOINC (bit 86) pushes no return address: the compiler uses it as a jump - // (the lowering treats it as one), so it never pairs with a RET. - const bool noinc_call = opcode == 0x944 && (instruction.words[2] & 0x400000) != 0; - if (opcode == 0x944 && !noinc_call) { + // CALL.REL.NOINC (bit 86) pushes no return address, but its callee still returns + // with RET.REL.NODEC through the link register set ahead of the call. + if (opcode == 0x944) { last_call_target = std::max(last_call_target, static_cast(target / 16)); } else { last_forward_target = std::max(last_forward_target, static_cast(target / 16)); @@ -38,7 +37,7 @@ class ProgramCaptureState { if (opcode == 0x949 || opcode == 0x958) computed_branch = true; if (computed_branch) return Result::Continue; if (opcode == 0x94d && instruction.Bits(12, 4) == 7 && index >= last_forward_target) { - if (!last_call_target) return Result::Complete; + if (index >= last_call_target) return Result::Complete; main_exited = true; } if (opcode == 0x950 && instruction.Bits(12, 4) == 7 && main_exited && diff --git a/src/shader_recompiler/frontend/maxwell/translate/impl/integer_shift_right.cpp b/src/shader_recompiler/frontend/maxwell/translate/impl/integer_shift_right.cpp index 2dccd2ca6d..f33d1fb68f 100644 --- a/src/shader_recompiler/frontend/maxwell/translate/impl/integer_shift_right.cpp +++ b/src/shader_recompiler/frontend/maxwell/translate/impl/integer_shift_right.cpp @@ -42,9 +42,10 @@ void SHR(TranslatorVisitor& v, u64 insn, const IR::U32& shift) { const IR::U32 zero{v.ir.Imm32(0)}; const IR::U32 safe_bits{v.ir.Imm32(32)}; - const IR::U1 is_negative{v.ir.ILessThan(result, zero, true)}; const IR::U1 is_safe{v.ir.ILessThan(shift, safe_bits, false)}; - const IR::U32 clamped_value{v.ir.Select(is_negative, v.ir.Imm32(-1), zero)}; + const IR::U32 clamped_value{shr.is_signed != 0 + ? IR::U32{v.ir.ShiftRightArithmetic(base, v.ir.Imm32(31))} + : zero}; result = IR::U32{v.ir.Select(is_safe, result, clamped_value)}; } v.X(shr.dest_reg, result); diff --git a/src/tests/video_core/ampere_lowering.cpp b/src/tests/video_core/ampere_lowering.cpp index 3d2ea62600..5b091ebdb4 100644 --- a/src/tests/video_core/ampere_lowering.cpp +++ b/src/tests/video_core/ampere_lowering.cpp @@ -147,4 +147,74 @@ TEST_CASE("Ampere programs go on past an EXIT that a branch jumps over", "[video REQUIRE(static_cast(lowered)); REQUIRE(lowered.source_size == sizeof(code)); } + +TEST_CASE("Ampere capture ends at the return of a helper called past the EXIT", + "[video_core][ampere]") { + // Ocarina 69a40579971fb867: MOV R0, 0x20; CALL.REL.NOINC; EXIT; RET.REL.NODEC R0, then the + // next program in memory. + const std::array code{{ + {{0x00007802, 0x20, 0x00000f00, 0}}, + {{0x00007944, 0x10, 0x03c00000, 0}}, + EXIT, + {{0x00007950, 0xffffffc0, 0x03c3ffff, 0}}, + }}; + ProgramCaptureState capture; + for (u32 index = 0; index + 1 < code.size(); ++index) + CHECK(capture.Observe(code[index], index) == ProgramCaptureState::Result::Continue); + CHECK(capture.Observe(code.back(), 3) == ProgramCaptureState::Result::Complete); + REQUIRE(static_cast(LowerToMaxwell(code))); + + ProgramCaptureState jump; + CHECK(jump.Observe({{0x00007944, 0x10, 0x03c00000, 0}}, 0) == ProgramCaptureState::Result::Continue); + CHECK(jump.Observe(EXIT, 1) == ProgramCaptureState::Result::Continue); + CHECK(jump.Observe(EXIT, 2) == ProgramCaptureState::Result::Complete); +} + +TEST_CASE("Ampere uniform select and bit operations reach every uniform register", + "[video_core][ampere]") { + // Ocarina e8decc3e5c0c184e: @UP0 USEL UR4, UR4, 0x20, UP1. UP0 and UP1 come from R239 into + // P6 and P5; the SEL runs under P6 and the guest predicates are restored. + const std::array select{{ + {{0x04040887, 0x20, 0x08800000, 0x000fe200}}, + EXIT, + }}; + const auto selected = LowerToMaxwell(select); + INFO(selected.error); + REQUIRE(static_cast(selected)); + const std::array expected_select{ + 0x010000000207f0fcULL, 0x38e800000607fffdULL, 0x38f000000607ef00ULL, + 0x5ca002800fc6dcdcULL, 0x38f000000607fd00ULL}; + CHECK(std::search(selected.instructions.begin(), selected.instructions.end(), + expected_select.begin(), expected_select.end()) != selected.instructions.end()); + + // Ocarina 2079791ecb288b83: UBREV UR4, UR17. + const std::array reverse{{ + {{0x000472be, 0x11, 0x08000000, 0x000fc800}}, + EXIT, + }}; + const auto reversed = LowerToMaxwell(reverse); + INFO(reversed.error); + REQUIRE(static_cast(reversed)); + CHECK(reversed.instructions.front() == 0x382801800007e9dcULL); + + // Ocarina 2079791ecb288b83: UFLO.U32.SH UR21, UR21. + const std::array leading{{ + {{0x001572bd, 0x15, 0x080e0400, 0x000fe400}}, + EXIT, + }}; + const auto found = LowerToMaxwell(leading); + INFO(found.error); + REQUIRE(static_cast(found)); + CHECK(found.instructions.front() == 0x5c3002000ed700edULL); + + // Ocarina 2079791ecb288b83: UPOPC UR10, UR17. + const std::array count{{ + {{0x000a72bf, 0x11, 0x08000000, 0x000fe200}}, + EXIT, + }}; + const auto counted = LowerToMaxwell(count); + INFO(counted.error); + REQUIRE(static_cast(counted)); + CHECK(counted.instructions.front() == 0x5c0800000e9700e2ULL); +} } // namespace From 502e6e2f24f9b8cc6ca9c6ff26555b8fb9cdf92e Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:24:26 +0100 Subject: [PATCH 04/11] shader: lower Ocarina guarded uniform compares and uniform SEL Ported from zortos293/Lichii 18e4efbbb8 (fix/oot-bloom). @UPn UISETP loads the guard and the old output into the predicate file, so a false guard merges the output back unchanged. SEL takes a uniform second source like FSEL and FMNMX. Assisted-by: Claude Opus 5.5 (ported the upstream patch hunks and tests, resolved conflicts) (cherry picked from commit edf4eacc7ddd816bbe5b90d780b571e07ac207a1) --- .../frontend/ampere/lower_to_maxwell.h | 21 ++++++++++----- src/tests/video_core/ampere_lowering.cpp | 27 +++++++++++++++++++ 2 files changed, 41 insertions(+), 7 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index fa90ec0642..e6a87ee56b 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -189,7 +189,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const u32 op = instruction.Bits(0, 12); uses_uniform_bank |= op == 0x4a4 || op == 0x2bf; uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0x896 || op == 0xf60 || op == 0xc11 || op == 0xd08 || 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 == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || op == 0x8a4 || op == 0xd05 || (op == 0x984 && instruction.Bits(91, 1)) || (op == 0x988 && instruction.Bits(91, 1)) || op == 0x89c || op == 0x883 || op == 0x2a5 || op == 0xd06 || op == 0xf63 || op == 0xf66 || op == 0xf69 || op == 0xc0b || op == 0x8a5 || op == 0x299 || op == 0xc09 || (op == 0x81c && instruction.Bits(67, 1)) || op == 0xc0a || + op == 0x886 || op == 0x3c2 || op == 0xc82 || op == 0x882 || op == 0x292 || op == 0x892 || op == 0x291 || op == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || op == 0x8a4 || op == 0xd05 || (op == 0x984 && instruction.Bits(91, 1)) || (op == 0x988 && instruction.Bits(91, 1)) || op == 0x89c || op == 0x883 || op == 0x2a5 || op == 0xd06 || op == 0xf63 || op == 0xf66 || op == 0xf69 || op == 0xc0b || op == 0x8a5 || op == 0x299 || op == 0xc09 || op == 0xc07 || (op == 0x81c && instruction.Bits(67, 1)) || op == 0xc0a || (op == 0x83b && instruction.Bits(91, 1)); } const auto address_of = [](u32 index) { return 0x50U + index * 8 + ((index + 1) / 3) * 8; }; @@ -1722,15 +1722,22 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const u32 set_op = insn.Bits(74, 2); const u32 a = uniform_register(insn.Bits(24, 8)), b = opcode == 0x88c ? 252U : uniform_register(insn.Bits(32, 8)); if (output > 6 || insn.Bits(84, 3) != 7 || - insn.Bits(72, 1) || set_op == 3 || insn.Bits(12, 4) != 7 || a == ~0U || b == ~0U) + insn.Bits(72, 1) || set_op == 3 || a == ~0U || b == ~0U) return fail(index, "native uniform comparison mode"); const u32 temporary = 6 - output; const u32 accum_pred = accum == 7 ? 7U : 6 - accum; - const u32 mask = (1U << temporary) | (accum == 7 ? 0U : (1U << accum_pred)); - if (opcode == 0x88c) mov_imm(252, insn.words[1]); + // @UPn: the guard and the old output load too, so a false guard merges it back unchanged. + const bool guarded = native_predicate != 7; + const u32 load_mask = (accum == 7 ? 0U : (1U << accum_pred)) | + (guarded ? (1U << temporary) | (1U << guard_slot) : 0U); + const u32 mask = (1U << temporary) | load_mask; + if (opcode == 0x88c) + emit(0x0100000000000000ULL | (u64{7} << 16) | 252 | (u64{15} << 12) | + (u64{insn.words[1]} << 20)); emit(0x38e8000000070000ULL | 253 | (u64{255} << 8) | (u64{mask} << 20)); - if (accum != 7) emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{1U << accum_pred} << 20)); - emit(0x5b60000000070000ULL | 7 | (u64{temporary} << 3) | + if (load_mask) emit(0x38f0000000070000ULL | (u64{239} << 8) | (u64{load_mask} << 20)); + emit(0x5b60000000000000ULL | + (u64{guarded ? (native_predicate & 8) | guard_slot : 7U} << 16) | 7 | (u64{temporary} << 3) | (u64{a} << 8) | (u64{b} << 20) | (u64{accum_pred} << 39) | (u64{accum_neg} << 42) | (u64{set_op} << 45) | (u64{insn.Bits(73, 1)} << 48) | (u64{insn.Bits(76, 3)} << 49)); @@ -2749,7 +2756,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{insn.Bits(78, 2)} << 39)); break; } - case 0x207: case 0x807: case 0xa07: + case 0x207: case 0x807: case 0xa07: case 0xc07: case 0x208: case 0x808: case 0xa08: { // Bitwise SEL / FSEL // FSEL applies the floating operand modifiers (negate/absolute) before selecting. const bool floating = (opcode & 0x1ff) == 0x008; diff --git a/src/tests/video_core/ampere_lowering.cpp b/src/tests/video_core/ampere_lowering.cpp index 5b091ebdb4..49fe914aea 100644 --- a/src/tests/video_core/ampere_lowering.cpp +++ b/src/tests/video_core/ampere_lowering.cpp @@ -217,4 +217,31 @@ TEST_CASE("Ampere uniform select and bit operations reach every uniform register REQUIRE(static_cast(counted)); CHECK(counted.instructions.front() == 0x5c0800000e9700e2ULL); } + +TEST_CASE("Ampere guarded uniform comparisons keep the output when the guard is false", + "[video_core][ampere]") { + // Ocarina 2079791ecb288b83: @UP4 UISETP.GT.U32.AND UP0, UPT, UR21, 0x1f, UPT. UP4 and the old + // UP0 load into P2 and P6, so a false guard merges UP0 back unchanged. + const std::array code{{ + {{0x1500488c, 0x1f, 0x0bf04070, 0x000fe200}}, + EXIT, + }}; + const auto lowered = LowerToMaxwell(code); + INFO(lowered.error); + REQUIRE(static_cast(lowered)); + const std::array expected{ + 0x0100000001f7f0fcULL, 0x38e800000447fffdULL, 0x38f000000447ef00ULL, + 0x5b6803800fc2ed37ULL, 0x38e800000407efefULL, 0x38f000000447fd00ULL}; + CHECK(std::equal(expected.begin(), expected.end(), lowered.instructions.begin())); + + // SEL R3, R4, UR7, P0 + const std::array select{{ + {{0x04037c07, 0x7, 0x08000000, 0x000fe400}}, + EXIT, + }}; + const auto selected = LowerToMaxwell(select); + INFO(selected.error); + REQUIRE(static_cast(selected)); + CHECK(selected.instructions.front() == 0x5ca000000df70403ULL); +} } // namespace From f9185051468175952d26c05c6fe1ae4a7cfe6749 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:33:31 +0100 Subject: [PATCH 05/11] shader: lower Ocarina 2D array, per-pixel offset and depth gathers TLD4 accepted only plain 2D gathers with an optional AOFFI. Ocarina also gathers from 2D arrays, with per-pixel offsets (bits 76-77 = 2) and with a depth reference (bit 78). The source lists follow Mesa NAK: Ra holds the coordinates with the array index first, Rb the offsets and then the reference, and with nothing for Rb the coordinates split evenly across both. Maxwell TLD4 reads the same lists, so the lowering copies them into its coordinate and meta registers and sets the type, PTP and DC fields. FMUL and FADD in register form get shadow operands for high registers. Assisted-by: Claude Opus 5.5 (decoded the rejections from the log, wrote the lowering and tests) (cherry picked from commit f1a5c6e04fbce7250b92e82d51291821a59f9f36) --- .../frontend/ampere/lower_to_maxwell.h | 45 +++++++++++------- src/tests/video_core/ampere_lowering.cpp | 46 +++++++++++++++++++ 2 files changed, 75 insertions(+), 16 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index e6a87ee56b..2097020e07 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -80,6 +80,7 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re case 0x805: operand(16, 2); break; case 0x308: operand(16); operand(32); break; case 0x433: case 0x633: operand(16); operand(24); break; + case 0x220: case 0x221: operand(16); operand(24); operand(32); break; case 0x233: operand(16); operand(24); operand(32); break; case 0xa20: case 0x621: case 0x808: case 0x923: operand(16); operand(24); break; case 0x209: operand(16); operand(24); operand(32); break; @@ -2251,31 +2252,43 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 } break; } - case 0xb63: case 0xf63: { // TLD4 (gather), 2D without offsets or depth compare; component at bits 87-88. + case 0xb63: case 0xf63: { // TLD4 (gather), 2D or 2D array; component at bits 87-88. // 0xf63: the handle in a uniform register (bit 91), Maxwell TLD4.B with the handle in Rb. const bool bindless = opcode == 0xf63; const u32 handle = bindless ? uniform_register(insn.Bits(40, 6)) : 255; const u32 mask = insn.Bits(72, 4), second_dest = insn.Bits(64, 8); const u32 source_a = insn.Bits(24, 8), source_b = insn.Bits(32, 8); const u32 binding = insn.Bits(54, 5); - // Bit 76: a texel offset after the coordinates (x, y in Ra, Ra+1, offsets in Rb, as - // for TEX); Maxwell TLD4 takes it as AOFFI (bits 54-55) from Rb. - const bool offsets = insn.Bits(76, 1) != 0; - if (insn.Bits(61, 3) != 1 || insn.Bits(77, 3) || insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || - (offsets && (bindless || source_b == 255 || (source_a != 255 && source_a >= 245))) || + // Dimension 5: 2D array. Bits 76-77: AOFFI (1, one packed offset) or per-pixel offsets + // (2, two words), bit 78: depth compare. As in Mesa NAK, Ra holds the coordinates (the + // array index first) and Rb the offsets and then the depth reference; with nothing for + // Rb the coordinate list splits evenly across them. Maxwell TLD4 takes the same lists. + const bool array = insn.Bits(61, 3) == 5; + const u32 offset_mode = insn.Bits(76, 2); + const bool depth = insn.Bits(78, 1) != 0; + const u32 coordinates = array ? 3U : 2U; + const u32 meta = (offset_mode == 1 ? 1U : offset_mode == 2 ? 2U : 0U) + (depth ? 1U : 0U); + const u32 in_a = meta ? coordinates : (coordinates + 1) / 2; + const u32 in_b = meta ? meta : coordinates - in_a; + if ((insn.Bits(61, 3) != 1 && !array) || offset_mode == 3 || insn.Bits(79, 1) || + insn.Bits(81, 3) != 7 || insn.Bits(59, 1) || + (bindless && (offset_mode || depth || array)) || + (offset_mode && source_b == 255) || + (source_a != 255 && source_a + in_a > 246) || (source_b != 255 && source_b + in_b > 246) || insn.Bits(89, 1) || mask == 0 || (bindless ? (handle >= 246 || insn.Bits(46, 13) || !insn.Bits(91, 1)) : (insn.Bits(40, 14) > 0x1fff || binding >= 18 || (result.required_texture_buffer != ~0U && result.required_texture_buffer != binding)))) return fail(index, "native texture gather mode"); if (!bindless) result.required_texture_buffer = binding; - // Two coordinates: x in Ra, y in Rb (the coordinate list splits evenly across them); - // with an offset, x and y in Ra, Ra+1 and the offset in Rb. - const u32 y = !offsets ? source_b : source_a == 255 ? 255U : source_a + 1; - emit(0x5c98000000000000ULL | predicate | 250 | (u64{15} << 39) | (u64{source_a} << 20)); - emit(0x5c98000000000000ULL | predicate | 251 | (u64{15} << 39) | (u64{y} << 20)); - if (offsets) { - emit(0x5c98000000000000ULL | predicate | 246 | (u64{15} << 39) | (u64{source_b} << 20)); + const auto element = [](u32 base, u32 offset) { return base == 255 ? 255U : base + offset; }; + for (u32 i = 0; i < coordinates; ++i) { + const u32 source = i < in_a ? element(source_a, i) : element(source_b, i - in_a); + emit(0x5c98000000000000ULL | predicate | (250 + i) | (u64{15} << 39) | (u64{source} << 20)); + } + for (u32 i = 0; i < meta; ++i) { + emit(0x5c98000000000000ULL | predicate | (246 + i) | (u64{15} << 39) | + (u64{element(source_b, i)} << 20)); } if (bindless) { emit(0x5c98000000000000ULL | predicate | 246 | (u64{15} << 39) | (u64{handle} << 20)); @@ -2284,9 +2297,9 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51)); } else { emit(0xc800000000000000ULL | predicate | 250 | (u64{250} << 8) | - (u64{offsets ? 246U : 255U} << 20) | (u64{2} << 28) | (u64{mask} << 31) | - (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{7} << 51) | - (u64{offsets ? 1U : 0U} << 54) | (u64{insn.Bits(87, 2)} << 56)); + (u64{meta ? 246U : 255U} << 20) | (u64{array ? 3U : 2U} << 28) | (u64{mask} << 31) | + (u64{insn.Bits(40, 14)} << 36) | (u64{insn.Bits(90, 1)} << 49) | (u64{depth} << 50) | + (u64{7} << 51) | (u64{offset_mode} << 54) | (u64{insn.Bits(87, 2)} << 56)); } const u32 components = static_cast(std::popcount(mask)); const bool split = second_dest != 255; diff --git a/src/tests/video_core/ampere_lowering.cpp b/src/tests/video_core/ampere_lowering.cpp index 49fe914aea..26825b92f7 100644 --- a/src/tests/video_core/ampere_lowering.cpp +++ b/src/tests/video_core/ampere_lowering.cpp @@ -244,4 +244,50 @@ TEST_CASE("Ampere guarded uniform comparisons keep the output when the guard is REQUIRE(static_cast(selected)); CHECK(selected.instructions.front() == 0x5ca000000df70403ULL); } + +TEST_CASE("Ampere gathers take arrays, per-pixel offsets and depth references", + "[video_core][ampere]") { + const auto lower = [](const Instruction& instruction) { + const std::array code{{instruction, EXIT}}; + return LowerToMaxwell(code); + }; + const auto field = [](u64 word, u32 offset, u32 count) { + return (word >> offset) & ((u64{1} << count) - 1); + }; + // Ocarina: TLD4.R.2D.ARRAY R44, R42, R8, R4. Layer and x in R8, R9, y in R4. + const auto array = lower({{0x082c7b63, 0xb0801004, 0x001e0f2a, 0x00014200}}); + INFO(array.error); + REQUIRE(static_cast(array)); + REQUIRE(array.instructions.size() >= 4); + CHECK(array.instructions[0] == 0x5c980780008700faULL); + CHECK(array.instructions[1] == 0x5c980780009700fbULL); + CHECK(array.instructions[2] == 0x5c980780004700fcULL); + CHECK(field(array.instructions[3], 20, 8) == 255); + CHECK(field(array.instructions[3], 28, 3) == 3); + CHECK(field(array.instructions[3], 50, 1) == 0); + CHECK(field(array.instructions[3], 54, 2) == 0); + + // Ocarina: @!P1 TLD4.R.PTP R24, R14, R6, R16: x, y in R6, R7, the offsets in R16, R17. + const auto per_pixel = lower({{0x06189b63, 0x30007c10, 0x001e2f0e, 0x000f4600}}); + INFO(per_pixel.error); + REQUIRE(static_cast(per_pixel)); + REQUIRE(per_pixel.instructions.size() >= 5); + CHECK(field(per_pixel.instructions[2], 20, 8) == 16); + CHECK(field(per_pixel.instructions[3], 20, 8) == 17); + CHECK(field(per_pixel.instructions[4], 20, 8) == 246); + CHECK(field(per_pixel.instructions[4], 28, 3) == 2); + CHECK(field(per_pixel.instructions[4], 54, 2) == 2); + + // Ocarina: TLD4.R.DC R12, R14, R24, R28: x, y in R24, R25, the reference in R28. + const auto depth = lower({{0x180c7b63, 0x30800e1c, 0x001e4f0e, 0x000f4200}}); + INFO(depth.error); + REQUIRE(static_cast(depth)); + REQUIRE(depth.instructions.size() >= 4); + CHECK(field(depth.instructions[2], 20, 8) == 28); + CHECK(field(depth.instructions[3], 50, 1) == 1); + CHECK(field(depth.instructions[3], 54, 2) == 0); + + // Offset mode 3 is not an offset form. + CHECK_FALSE(lower({{0x082c7b63, 0xb0801004, 0x001e3f2a, 0x00014200}})); +} } // namespace From 8424eeb2531a3bfb4788918405996f21adf17365 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 01:46:19 +0100 Subject: [PATCH 06/11] shader: lower the rest of Ocarina's dumped native programs All 833 programs from an Ocarina shader dump now lower; the 800 that already did lower to identical code. - SGXT takes its signed form (bit 73) as BFE.S32. - USHF.L.U32.HI with an immediate (0x499) lowers like SHF 0x419, with a UPn guard; Ocarina's mesh programs use it. - UIADD3 carries go to and come from any UPn, through P(6-n) and R239 as UISETP does, not only UP0 and UP1. - LEA.HI takes a negated Rb (bit 63) when it is not .X. - SURED takes every op (bits 87-89), which Maxwell SURED reads in the same order at bits 24-26. - Shadow operands for LDG, IADD3, FFMA, FMUL, FADD, FSEL, IMNMX, SGXT, S2R, ISETP and I2F forms the dump uses. An operand may no longer straddle the reserved boundary or span two borrowed quartets; either would be renamed to registers that are not adjacent. Assisted-by: Claude Opus 5.5 (lowered the dump offline, decoded the rejections, wrote the fixes and tests) (cherry picked from commit 19e25d0cfb7ec570f6e613fcee92dbdcdc6dc873) --- .../frontend/ampere/lower_to_maxwell.h | 63 +++++++++++++----- src/tests/video_core/ampere_lowering.cpp | 66 +++++++++++++++++++ 2 files changed, 113 insertions(+), 16 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index 2097020e07..c461c56b28 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -81,6 +81,15 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re case 0x308: operand(16); operand(32); break; case 0x433: case 0x633: operand(16); operand(24); break; case 0x220: case 0x221: operand(16); operand(24); operand(32); break; + case 0x421: case 0x817: case 0x820: operand(16); operand(24); break; + case 0x81a: operand(16); operand(24); break; + case 0x919: operand(16); break; + case 0x20c: operand(24); operand(32); break; + case 0x208: operand(16); operand(24); operand(32); break; + case 0x823: operand(16); operand(24); operand(64); break; + case 0x245: operand(16); operand(32); break; + case 0x810: operand(16); operand(24); operand(64); break; + case 0x210: case 0x223: operand(16); operand(24); operand(32); operand(64); break; case 0x233: operand(16); operand(24); operand(32); break; case 0xa20: case 0x621: case 0x808: case 0x923: operand(16); operand(24); break; case 0x209: operand(16); operand(24); operand(32); break; @@ -119,6 +128,10 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re operand(opcode == 0x981 ? 16 : 32, count); operand(24, input.Bits(72, 1) ? 2 : 1); break; } + case 0x381: { // LDG with a 64-bit address + const u32 width = input.Bits(73, 3), count = width == 6 ? 4 : width == 5 ? 2 : 1; + operand(16, count); operand(24, 2); break; + } default: result.valid = false; return result; } @@ -127,7 +140,12 @@ inline RegisterPressure PrepareRegisterPressure(const Instruction& input, u32 re for (const auto [bit, count] : operands) { const u32 reg = input.Bits(bit, 8); if (reg == 255) continue; - if (reg + count > 255) { result.valid = false; return result; } + // Each quartet is borrowed on its own, so an operand above `reserved` must stay in one, + // and none may straddle `reserved`. + if (reg + count > 255 || (reg < reserved && reg + count > reserved) || + (reg >= reserved && reg / 4 != (reg + count - 1) / 4)) { + result.valid = false; return result; + } for (u32 component = 0; component < count; ++component) { used[reg + component] = true; if (reg + component >= reserved) groups[(reg + component) / 4] = true; @@ -188,7 +206,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 }); for (const auto& instruction : code) { const u32 op = instruction.Bits(0, 12); - uses_uniform_bank |= op == 0x4a4 || op == 0x2bf; + uses_uniform_bank |= op == 0x4a4 || op == 0x2bf || op == 0x499; uses_uniform_bank |= op == 0x3c4 || op == 0x2bd || op == 0x896 || op == 0xf60 || op == 0xc11 || op == 0xd08 || 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 == 0x981 || op == 0xc13 || op == 0xc10 || op == 0xc23 || op == 0x526 || op == 0x2be || op == 0x8a4 || op == 0xd05 || (op == 0x984 && instruction.Bits(91, 1)) || (op == 0x988 && instruction.Bits(91, 1)) || op == 0x89c || op == 0x883 || op == 0x2a5 || op == 0xd06 || op == 0xf63 || op == 0xf66 || op == 0xf69 || op == 0xc0b || op == 0x8a5 || op == 0x299 || op == 0xc09 || op == 0xc07 || (op == 0x81c && instruction.Bits(67, 1)) || op == 0xc0a || (op == 0x83b && instruction.Bits(91, 1)); @@ -328,7 +346,7 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 } // 0x899 (USHF) lowers to one SHR/SHL, so a uniform guard is swapped in like these. const bool uniform_predication = opcode == 0xabb || opcode == 0xab9 || opcode == 0xc82 || opcode == 0x882 || - opcode == 0x899; + opcode == 0x899 || opcode == 0x499; const u32 native_predicate = insn.Bits(12, 4); const bool guarded_uniform = (uniform_predication || uniform_guard) && native_predicate != 7; // A uniform guard UPn lives in R239 bit 6-n and is swapped into P(6-n) for the op. @@ -1379,19 +1397,19 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 if (strong) emit(0xef98000000000000ULL | predicate | (u64{1} << 8)); break; } - case 0x59f: { // SURED.D.2D.ADD.STRONG.GPU.IGN U32, direct descriptor (title 1's race compute). + case 0x59f: { // SURED.D.2D.{ADD..XOR}.STRONG.GPU.IGN U32, direct descriptor (op in bits 87-89). const u32 binding = insn.Bits(54, 5), offset = insn.Bits(40, 14); const u32 data = insn.Bits(32, 8), coordinate = insn.Bits(24, 8); - if (insn.words[2] != 0x0010e000 || insn.Bits(61, 3) != 3 || binding >= 18 || + if ((insn.words[2] & ~0x03800000U) != 0x0010e000 || insn.Bits(61, 3) != 3 || binding >= 18 || offset > 0x1fff || insn.Bits(59, 2) || coordinate > 253) return fail(index, "native surface reduction 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; - // Maxwell SURED (as the frontend decodes it): operand 0-7, coordinates 8-15, ADD and - // U32 are zero at 24 and 20, type 33 (3 = 2D), bound offset 36, bound 51. + // Maxwell SURED (as the frontend decodes it): operand 0-7, coordinates 8-15, op 24 (the + // same ADD..XOR order), U32 zero at 20, type 33 (3 = 2D), bound offset 36, bound 51. emit(0xeb40000000000000ULL | predicate | data | (u64{coordinate} << 8) | - (u64{3} << 33) | (u64{offset} << 36) | (u64{1} << 51)); + (u64{insn.Bits(87, 3)} << 24) | (u64{3} << 33) | (u64{offset} << 36) | (u64{1} << 51)); break; } case 0x301: // BREV register. @@ -2696,11 +2714,11 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{insn.Bits(76, 2)} << 52) | (u64{insn.Bits(78, 1)} << 55)); break; } - case 0x81a: { // SGXT.U32 Rd, Ra, n: keep the low n bits (BFE.U32 at position 0) - if (insn.words[2] || insn.words[1] > 32 || dest >= 255) + case 0x81a: { // SGXT[.U32] Rd, Ra, n: the low n bits, sign-extended unless .U32 (BFE at position 0) + if ((insn.words[2] & ~0x200U) || insn.words[1] > 32 || dest >= 255) return fail(index, "native sign extend mode"); emit(base(0x3800000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | - (u64{insn.words[1] << 8} << 20)); + (u64{insn.words[1] << 8} << 20) | (u64{insn.Bits(73, 1)} << 48)); break; } case 0x299: { // USHF by a uniform amount, one register shifted (SHL/SHR, bit 75 wraps). @@ -2719,6 +2737,17 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{source} << 8) | (u64{shift} << 20) | (u64{insn.Bits(75, 1)} << 39)); break; } + case 0x499: { // USHF.L[.W].U32.HI URd, URZ, URb, imm: imm << URb, as 0x419 below + const u32 target = uniform_register(static_cast(dest)); + const u32 shift = uniform_register(insn.Bits(64, 8)); + if (insn.Bits(24, 8) != 63 || (insn.words[2] & ~0x8ffU) != 0x08010600 || + (insn.words[3] & 0xc00001ffU) || target == ~0U || target == 255 || shift == ~0U) + return fail(index, "native uniform immediate high-word shift mode"); + mov_imm(250, insn.words[1]); + emit(0x5c48000000000000ULL | predicate | target | (u64{250} << 8) | + (u64{shift} << 20) | (u64{insn.Bits(75, 1)} << 39)); + break; + } case 0x419: { // SHF.L[.W].U32.HI Rd, RZ, Rb, imm (high word of imm:0 << Rb) = imm << Rb // .W (bit 75) wraps the shift to Rb & 31; without it a shift of 32 or more gives zero, // like Maxwell SHL without .W (bit 39). @@ -3006,7 +3035,8 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const bool extended = insn.Bits(74, 1) != 0; const u32 high = insn.Bits(64, 8); if (insn.Bits(72, 1) || insn.Bits(81, 3) != 7 || - (high != 255 && insn.Bits(73, 1)) || (opcode == 0x211 && insn.Bits(62, 2)) || + (high != 255 && insn.Bits(73, 1)) || + (opcode == 0x211 && (insn.Bits(62, 1) || (insn.Bits(63, 1) && extended))) || (opcode == 0xc11 && (insn.words[1] & ~0x3fU))) return fail(index, "native high effective address mode"); u32 b = opcode == 0xc11 ? uniform_register(insn.Bits(32, 8)) : insn.Bits(32, 8); @@ -3032,7 +3062,9 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 (u64{insn.Bits(87, 3)} << 39) | (u64{insn.Bits(90, 1)} << 42)); emit(0x5c10000000000000ULL | predicate | 250 | (u64{250} << 8) | (u64{252} << 20)); } - emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{b} << 20)); + // 0x211 bit 63: -Rb, Maxwell IADD's negated second source (bit 48). + emit(base(0x5c10000000000000ULL) | (u64{250} << 8) | (u64{b} << 20) | + (u64{opcode == 0x211 ? insn.Bits(63, 1) : 0U} << 48)); } break; } @@ -3077,9 +3109,8 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const bool carry_in = insn.Bits(74, 1) != 0; const u32 carry_out = insn.Bits(81, 3); const u32 a = insn.Bits(24, 8); - if ((carry_in && carry_out != 7) || (opcode == 0x210 && (insn.words[1] & ~0xffU)) || - (rewritten_uniform && carry_out != 7 && carry_out > 1) || - (rewritten_uniform && carry_in && insn.Bits(87, 3) > 1 && insn.Bits(87, 3) != 7)) + // A uniform carry goes through P(6-n) and R239 bit 6-n, so any UPn works. + if ((carry_in && carry_out != 7) || (opcode == 0x210 && (insn.words[1] & ~0xffU))) return fail(index, "native integer carry mode"); if (opcode == 0x810) { emit(0x1c00000000000000ULL | predicate | 250 | (u64{a} << 8) | (u64{insn.words[1]} << 20)); diff --git a/src/tests/video_core/ampere_lowering.cpp b/src/tests/video_core/ampere_lowering.cpp index 26825b92f7..0ea19ad9a6 100644 --- a/src/tests/video_core/ampere_lowering.cpp +++ b/src/tests/video_core/ampere_lowering.cpp @@ -290,4 +290,70 @@ TEST_CASE("Ampere gathers take arrays, per-pixel offsets and depth references", // Offset mode 3 is not an offset form. CHECK_FALSE(lower({{0x082c7b63, 0xb0801004, 0x001e3f2a, 0x00014200}})); } + +TEST_CASE("Ampere Ocarina sign extends, uniform shifts and carries, LEA.HI and SURED forms", + "[video_core][ampere]") { + const auto lower = [](const Instruction& instruction) { + const std::array code{{instruction, EXIT}}; + return LowerToMaxwell(code); + }; + const auto field = [](u64 word, u32 offset, u32 count) { + return (word >> offset) & ((u64{1} << count) - 1); + }; + // SGXT R32, R32, 0x1: signed, so BFE.S32 (bit 48). + const auto extend = lower({{0x2020781a, 1, 0x00000200, 0x000fe400}}); + INFO(extend.error); + REQUIRE(static_cast(extend)); + CHECK(extend.instructions.front() == 0x3801000010072020ULL); + + // USHF.L.U32.HI UR5, URZ, UR4, 0x1: 1 << UR4 into UR5 (R221). + const auto shift = lower({{0x3f057499, 1, 0x08010e04, 0x000fe400}}); + INFO(shift.error); + REQUIRE(static_cast(shift)); + REQUIRE(shift.instructions.size() >= 2); + CHECK(shift.instructions[1] == 0x5c4800800dc7faddULL); + + // UIADD3 UR6, UP2, UR15, UR6, URZ: the carry goes through P4 into R239 bit 4. + const auto carry = lower({{0x0f067290, 6, 0x0ff5e03f, 0x000fe200}}); + INFO(carry.error); + REQUIRE(static_cast(carry)); + const std::array expected_carry{ + 0x5c1000000de7e7faULL, 0x38e800000107fffdULL, 0x5b6203800e77fa27ULL, + 0x38e800000107efefULL, 0x38f000000107fd00ULL, 0x5c9807800fa700deULL}; + CHECK(std::equal(expected_carry.begin(), expected_carry.end(), carry.instructions.begin())); + + // LEA.HI R1, R1, -R0, RZ, 0x1e: the final add negates R0. + const auto lea = lower({{0x01017211, 0x80000000, 0x078ff0ff, 0x000fca00}}); + INFO(lea.error); + REQUIRE(static_cast(lea)); + REQUIRE(lea.instructions.size() >= 2); + CHECK(lea.instructions[1] == 0x5c1100000007fa01ULL); + + // SURED.OR: op 6 at Maxwell bits 24-26. + const auto reduce = lower({{0x0400159f, 0x60008400, 0x0310e000, 0x0001e200}}); + INFO(reduce.error); + REQUIRE(static_cast(reduce)); + CHECK(field(reduce.instructions.front(), 24, 3) == 6); + + // LDG.E.128 R216, [R216.64] in a program using uniform registers: R216-R219 are reserved, + // so the load goes through a borrowed low quartet. + const std::array load{{ + {{0x00007c02, 0, 0x08000f00, 0}}, + {{0xd8d87381, 0, 0x001ead00, 0x000f2800}}, + EXIT, + }}; + const auto loaded = LowerToMaxwell(load); + INFO(loaded.error); + REQUIRE(static_cast(loaded)); + CHECK(std::any_of(loaded.instructions.begin(), loaded.instructions.end(), [&](u64 word) { + return (word >> 52) == 0xeed && field(word, 0, 8) < 216 && field(word, 8, 8) < 216; + })); + // A reserved operand spanning two quartets (R218-R221) cannot be borrowed for. + const std::array straddle{{ + {{0x00007c02, 0, 0x08000f00, 0}}, + {{0xd8da7381, 0, 0x001ead00, 0x000f2800}}, + EXIT, + }}; + CHECK_FALSE(LowerToMaxwell(straddle)); +} } // namespace From 01783ceb29d2d94bda72f273003323a2a446c7ae Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 02:00:35 +0100 Subject: [PATCH 07/11] shader: route untracked 32-bit global atomics to their storage buffer Every SPIR-V global atomic threw "SPIR-V Instruction is not implemented", so a compute program whose atomic address the storage buffer pass could not track failed to build and its dispatches were dropped (four Ocarina compute pipelines; the dumps show U32 ADD/MAX, S32 MIN and F32 ADD reductions). Each 32-bit global atomic the program uses now gets a function on the routing global loads and stores already use: the address picks the candidate storage buffer, and the operation runs on its word. INC, DEC and F32 add go through the storage buffer CAS loops. 64-bit and packed forms still throw. The test emits each of the 12 operations untracked; the modules pass spirv-val (Vulkan 1.1, standard uniform buffer layout). Assisted-by: Claude Opus 5.5 (traced the pipeline failures, wrote the routing and test) (cherry picked from commit 6b934346d76cf5a6753c9279df22f9df830fade4) --- .../backend/spirv/emit_spirv_atomic.cpp | 56 ++++---- .../backend/spirv/emit_spirv_instructions.h | 24 ++-- .../backend/spirv/spirv_emit_context.cpp | 83 ++++++++++- .../backend/spirv/spirv_emit_context.h | 7 +- src/tests/CMakeLists.txt | 1 + .../video_core/global_atomic_routing.cpp | 131 ++++++++++++++++++ 6 files changed, 260 insertions(+), 42 deletions(-) create mode 100644 src/tests/video_core/global_atomic_routing.cpp diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp index d1c08c1131..704c886df4 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp @@ -107,6 +107,14 @@ Id StorageAtomicU32x2(EmitContext& ctx, const IR::Value& binding, const IR::Valu ctx.OpStore(pointer, result); return original_value; } + +Id GlobalAtomic(EmitContext& ctx, EmitContext::GlobalAtomic kind, Id type, Id address, Id value) { + const Id func{ctx.global_atomic_funcs[static_cast(kind)]}; + if (!Sirit::ValidId(func)) { + throw NotImplementedException("SPIR-V Instruction"); + } + return ctx.OpFunctionCall(type, func, address, value); +} } // Anonymous namespace Id EmitSharedAtomicIAdd32(EmitContext& ctx, Id offset, Id value) { @@ -417,48 +425,48 @@ Id EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const I return ctx.OpPackHalf2x16(ctx.U32[1], result); } -Id EmitGlobalAtomicIAdd32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicIAdd32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::IAdd, ctx.U32[1], address, value); } -Id EmitGlobalAtomicSMin32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicSMin32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::SMin, ctx.U32[1], address, value); } -Id EmitGlobalAtomicUMin32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicUMin32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::UMin, ctx.U32[1], address, value); } -Id EmitGlobalAtomicSMax32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicSMax32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::SMax, ctx.U32[1], address, value); } -Id EmitGlobalAtomicUMax32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicUMax32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::UMax, ctx.U32[1], address, value); } -Id EmitGlobalAtomicInc32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicInc32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::Inc, ctx.U32[1], address, value); } -Id EmitGlobalAtomicDec32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicDec32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::Dec, ctx.U32[1], address, value); } -Id EmitGlobalAtomicAnd32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicAnd32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::And, ctx.U32[1], address, value); } -Id EmitGlobalAtomicOr32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicOr32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::Or, ctx.U32[1], address, value); } -Id EmitGlobalAtomicXor32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicXor32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::Xor, ctx.U32[1], address, value); } -Id EmitGlobalAtomicExchange32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicExchange32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::Exchange, ctx.U32[1], address, value); } Id EmitGlobalAtomicIAdd64(EmitContext&) { @@ -549,8 +557,8 @@ Id EmitGlobalAtomicExchange32x2(EmitContext&) { throw NotImplementedException("SPIR-V Instruction"); } -Id EmitGlobalAtomicAddF32(EmitContext&) { - throw NotImplementedException("SPIR-V Instruction"); +Id EmitGlobalAtomicAddF32(EmitContext& ctx, Id address, Id value) { + return GlobalAtomic(ctx, EmitContext::GlobalAtomic::AddF32, ctx.F32[1], address, value); } Id EmitGlobalAtomicAddF16x2(EmitContext&) { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h index b1d8ca707a..325dc22a2f 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h +++ b/src/shader_recompiler/backend/spirv/emit_spirv_instructions.h @@ -422,17 +422,17 @@ Id EmitStorageAtomicMaxF16x2(EmitContext& ctx, const IR::Value& binding, const I Id value); Id EmitStorageAtomicMaxF32x2(EmitContext& ctx, const IR::Value& binding, const IR::Value& offset, Id value); -Id EmitGlobalAtomicIAdd32(EmitContext& ctx); -Id EmitGlobalAtomicSMin32(EmitContext& ctx); -Id EmitGlobalAtomicUMin32(EmitContext& ctx); -Id EmitGlobalAtomicSMax32(EmitContext& ctx); -Id EmitGlobalAtomicUMax32(EmitContext& ctx); -Id EmitGlobalAtomicInc32(EmitContext& ctx); -Id EmitGlobalAtomicDec32(EmitContext& ctx); -Id EmitGlobalAtomicAnd32(EmitContext& ctx); -Id EmitGlobalAtomicOr32(EmitContext& ctx); -Id EmitGlobalAtomicXor32(EmitContext& ctx); -Id EmitGlobalAtomicExchange32(EmitContext& ctx); +Id EmitGlobalAtomicIAdd32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicSMin32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicUMin32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicSMax32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicUMax32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicInc32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicDec32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicAnd32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicOr32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicXor32(EmitContext& ctx, Id address, Id value); +Id EmitGlobalAtomicExchange32(EmitContext& ctx, Id address, Id value); Id EmitGlobalAtomicIAdd64(EmitContext& ctx); Id EmitGlobalAtomicSMin64(EmitContext& ctx); Id EmitGlobalAtomicUMin64(EmitContext& ctx); @@ -455,7 +455,7 @@ Id EmitGlobalAtomicAnd32x2(EmitContext& ctx); Id EmitGlobalAtomicOr32x2(EmitContext& ctx); Id EmitGlobalAtomicXor32x2(EmitContext& ctx); Id EmitGlobalAtomicExchange32x2(EmitContext& ctx); -Id EmitGlobalAtomicAddF32(EmitContext& ctx); +Id EmitGlobalAtomicAddF32(EmitContext& ctx, Id address, Id value); Id EmitGlobalAtomicAddF16x2(EmitContext& ctx); Id EmitGlobalAtomicAddF32x2(EmitContext& ctx); Id EmitGlobalAtomicMinF16x2(EmitContext& ctx); diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index f8f966a6b5..2b8f2aaf83 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -8,6 +8,7 @@ #include #include #include +#include #include @@ -493,7 +494,7 @@ EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_inf DefineImages(program.info, image_binding, bindings.image_scaling_index); DefineAttributeMemAccess(program.info); DefineWriteStorageCasLoopFunction(program.info); - DefineGlobalMemoryFunctions(program.info); + DefineGlobalMemoryFunctions(program); DefineRescalingInput(program.info); DefineRenderArea(program.info); } @@ -948,7 +949,8 @@ void EmitContext::DefineWriteStorageCasLoopFunction(const Info& info) { write_storage_cas_loop_func = func; } -void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { +void EmitContext::DefineGlobalMemoryFunctions(const IR::Program& program) { + const Info& info{program.info}; if (!info.uses_global_memory || !profile.support_int64) { return; } @@ -977,7 +979,7 @@ void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { const Id ssbo_offset{OpUConvert(U32[1], OpISub(U64, addr, ssbo_addr))}; const Id ssbo_index{OpShiftRightLogical(U32[1], ssbo_offset, Const(shift))}; const Id ssbo_pointer{OpAccessChain(element_pointer, ssbo_id, zero, ssbo_index)}; - callback(ssbo_pointer); + callback(ssbo_pointer, ssbo_id, ssbo_index); AddLabel(else_label); }}; if (info.native_global_bounds) { @@ -1021,7 +1023,7 @@ void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { const Id func_id{OpFunction(type, spv::FunctionControlMask::MaskNone, function_type)}; const Id addr{OpFunctionParameter(U64)}; define_body(ssbo_member, addr, element_pointer, shift, - [&](Id ssbo_pointer) { OpReturnValue(OpLoad(type, ssbo_pointer)); }); + [&](Id ssbo_pointer, Id, Id) { OpReturnValue(OpLoad(type, ssbo_pointer)); }); OpReturnValue(ConstantNull(type)); OpFunctionEnd(); return func_id; @@ -1031,7 +1033,7 @@ void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { const Id func_id{OpFunction(void_id, spv::FunctionControlMask::MaskNone, function_type)}; const Id addr{OpFunctionParameter(U64)}; const Id data{OpFunctionParameter(type)}; - define_body(ssbo_member, addr, element_pointer, shift, [&](Id ssbo_pointer) { + define_body(ssbo_member, addr, element_pointer, shift, [&](Id ssbo_pointer, Id, Id) { OpStore(ssbo_pointer, data); OpReturn(); }); @@ -1053,6 +1055,77 @@ void EmitContext::DefineGlobalMemoryFunctions(const Info& info) { define(&StorageDefinitions::U32x2, storage_types.U32x2, U32[2], sizeof(u32[2])); std::tie(load_global_func_u32x4, write_global_func_u32x4) = define(&StorageDefinitions::U32x4, storage_types.U32x4, U32[4], sizeof(u32[4])); + + // Atomics the storage buffer pass could not track: the same routing, with the operation on + // the routed word (INC, DEC and F32 add through their storage buffer CAS loops). + using enum GlobalAtomic; + const auto kind_of = [](IR::Opcode opcode) -> std::optional { + switch (opcode) { + case IR::Opcode::GlobalAtomicIAdd32: return IAdd; + case IR::Opcode::GlobalAtomicSMin32: return SMin; + case IR::Opcode::GlobalAtomicUMin32: return UMin; + case IR::Opcode::GlobalAtomicSMax32: return SMax; + case IR::Opcode::GlobalAtomicUMax32: return UMax; + case IR::Opcode::GlobalAtomicInc32: return Inc; + case IR::Opcode::GlobalAtomicDec32: return Dec; + case IR::Opcode::GlobalAtomicAnd32: return And; + case IR::Opcode::GlobalAtomicOr32: return Or; + case IR::Opcode::GlobalAtomicXor32: return Xor; + case IR::Opcode::GlobalAtomicExchange32: return Exchange; + case IR::Opcode::GlobalAtomicAddF32: return AddF32; + default: return std::nullopt; + } + }; + std::array(Count)> used{}; + for (const IR::Block* const block : program.blocks) { + for (const IR::Inst& inst : block->Instructions()) { + if (const auto kind = kind_of(inst.GetOpcode())) { + used[static_cast(*kind)] = true; + } + } + } + const Id scope{Const(static_cast(spv::Scope::Device))}; + for (size_t index = 0; index < used.size(); ++index) { + if (!used[index]) { + continue; + } + const auto kind = static_cast(index); + const Id type{kind == AddF32 ? F32[1] : U32[1]}; + const Id func_id{OpFunction(type, spv::FunctionControlMask::MaskNone, + TypeFunction(type, U64, type))}; + const Id addr{OpFunctionParameter(U64)}; + const Id value{OpFunctionParameter(type)}; + define_body(&StorageDefinitions::U32, addr, storage_types.U32.element, 2, + [&](Id pointer, Id ssbo, Id word) { + const auto atomic = [&](Id (Sirit::Module::*op)(Id, Id, Id, Id, Id)) { + return (this->*op)(U32[1], pointer, scope, zero, value); + }; + switch (kind) { + case IAdd: OpReturnValue(atomic(&Sirit::Module::OpAtomicIAdd)); break; + case SMin: OpReturnValue(atomic(&Sirit::Module::OpAtomicSMin)); break; + case UMin: OpReturnValue(atomic(&Sirit::Module::OpAtomicUMin)); break; + case SMax: OpReturnValue(atomic(&Sirit::Module::OpAtomicSMax)); break; + case UMax: OpReturnValue(atomic(&Sirit::Module::OpAtomicUMax)); break; + case And: OpReturnValue(atomic(&Sirit::Module::OpAtomicAnd)); break; + case Or: OpReturnValue(atomic(&Sirit::Module::OpAtomicOr)); break; + case Xor: OpReturnValue(atomic(&Sirit::Module::OpAtomicXor)); break; + case Exchange: OpReturnValue(atomic(&Sirit::Module::OpAtomicExchange)); break; + case Inc: + OpReturnValue(OpFunctionCall(U32[1], increment_cas_ssbo, word, value, ssbo)); + break; + case Dec: + OpReturnValue(OpFunctionCall(U32[1], decrement_cas_ssbo, word, value, ssbo)); + break; + case AddF32: + OpReturnValue(OpFunctionCall(F32[1], f32_add_cas, word, value, ssbo)); + break; + case Count: break; + } + }); + OpReturnValue(ConstantNull(type)); + OpFunctionEnd(); + global_atomic_funcs[index] = func_id; + } } void EmitContext::DefineRescalingInput(const Info& info) { diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h index 1f9b534aaf..a47a903ae7 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -343,6 +343,11 @@ class EmitContext final : public Sirit::Module { Id write_global_func_u32{}; Id write_global_func_u32x2{}; Id write_global_func_u32x4{}; + // 32-bit global atomics on untracked addresses, routed to the storage buffer holding them. + enum class GlobalAtomic : size_t { + IAdd, SMin, UMin, SMax, UMax, Inc, Dec, And, Or, Xor, Exchange, AddF32, Count, + }; + std::array(GlobalAtomic::Count)> global_atomic_funcs{}; bool need_input_position_indirect{}; Id input_position{}; @@ -420,7 +425,7 @@ class EmitContext final : public Sirit::Module { void DefineImages(const Info& info, u32& binding, u32& scaling_index); void DefineAttributeMemAccess(const Info& info); void DefineWriteStorageCasLoopFunction(const Info& info); - void DefineGlobalMemoryFunctions(const Info& info); + void DefineGlobalMemoryFunctions(const IR::Program& program); void DefineRescalingInput(const Info& info); void DefineRescalingInputPushConstant(); void DefineRescalingInputUniformConstant(); diff --git a/src/tests/CMakeLists.txt b/src/tests/CMakeLists.txt index 998a62821a..8e73d77108 100644 --- a/src/tests/CMakeLists.txt +++ b/src/tests/CMakeLists.txt @@ -40,6 +40,7 @@ add_executable(tests video_core/shader_invalidation.cpp video_core/descriptor_table.cpp video_core/global_address_tracking.cpp + video_core/global_atomic_routing.cpp video_core/nvn_storage_descriptor.cpp video_core/ampere_f2i.cpp video_core/ampere_ldc.cpp diff --git a/src/tests/video_core/global_atomic_routing.cpp b/src/tests/video_core/global_atomic_routing.cpp new file mode 100644 index 0000000000..5bce98ddd6 --- /dev/null +++ b/src/tests/video_core/global_atomic_routing.cpp @@ -0,0 +1,131 @@ +// SPDX-License-Identifier: GPL-3.0-or-later + +#include +#include + +#include + +#include "shader_recompiler/backend/spirv/emit_spirv.h" +#include "shader_recompiler/environment.h" +#include "shader_recompiler/frontend/ir/basic_block.h" +#include "shader_recompiler/frontend/ir/ir_emitter.h" +#include "shader_recompiler/frontend/ir/program.h" +#include "shader_recompiler/host_translate_info.h" +#include "shader_recompiler/ir_opt/passes.h" +#include "shader_recompiler/profile.h" + +namespace { +using namespace Shader; + +class StubEnvironment final : public Environment { +public: + StubEnvironment() { + stage = Stage::Compute; + } + u64 ReadInstruction(u32) override { return 0; } + u32 ReadCbufValue(u32, u32) override { return 0; } + TextureType ReadTextureType(u32) override { return {}; } + TexturePixelFormat ReadTexturePixelFormat(u32) override { return {}; } + bool IsTexturePixelFormatInteger(u32) override { return false; } + u32 ReadViewportTransformState() override { return 0; } + u32 TextureBoundBuffer() const override { return 0; } + u32 LocalMemorySize() const override { return 0; } + u32 SharedMemorySize() const override { return 0; } + std::array WorkgroupSize() const override { return {1, 1, 1}; } + bool HasHLEMacroState() const override { return false; } + std::optional GetReplaceConstBuffer(u32, u32) override { + return std::nullopt; + } + void Dump(u64, u64) override {} +}; + +// SPIR-V instructions with opcode `opcode`, as their operand words. +std::vector> Instructions(const std::vector& code, u32 opcode) { + std::vector> result; + for (size_t i = 5; i < code.size();) { + const u32 count = code[i] >> 16; + REQUIRE(count > 0); + REQUIRE(i + count <= code.size()); + if ((code[i] & 0xffff) == opcode) { + result.emplace_back(code.begin() + i + 1, code.begin() + i + count); + } + i += count; + } + return result; +} + +// A compute program whose atomic address comes from global memory, so the storage buffer pass +// tracks the load but not the atomic. +std::vector EmitUntrackedAtomic(IR::Opcode opcode) { + ObjectPool pool; + IR::Block block{pool}; + IR::IREmitter ir{block}; + const auto low = ir.GetCbuf(ir.Imm32(0), ir.Imm32(0x110)); + const auto high = ir.GetCbuf(ir.Imm32(0), ir.Imm32(0x114)); + const auto loaded = ir.LoadGlobal64(ir.PackUint2x32(ir.CompositeConstruct(low, high))); + const IR::U64 address{ir.PackUint2x32( + ir.CompositeConstruct(ir.CompositeExtract(loaded, 0), ir.CompositeExtract(loaded, 1)))}; + const IR::Value value{opcode == IR::Opcode::GlobalAtomicAddF32 ? IR::Value{ir.Imm32(1.0f)} + : IR::Value{ir.Imm32(1)}}; + const IR::Value previous{&*block.PrependNewInst(block.end(), opcode, {address, value})}; + const IR::U32 result{opcode == IR::Opcode::GlobalAtomicAddF32 + ? IR::U32{ir.BitCast(IR::F32{previous})} + : IR::U32{previous}}; + ir.WriteGlobal32(address, result); + + IR::Program program; + program.stage = Stage::Compute; + program.workgroup_size = {1, 1, 1}; + program.blocks = {&block}; + program.post_order_blocks = {&block}; + program.syntax_list = {{{.block = &block}, IR::AbstractSyntaxNode::Type::Block}, + {{}, IR::AbstractSyntaxNode::Type::Return}}; + HostTranslateInfo host{}; + host.min_ssbo_alignment = 16; + StubEnvironment env; + Optimization::GlobalMemoryToStorageBufferPass(env, program, host); + REQUIRE(std::ranges::any_of(block.Instructions(), [&](const IR::Inst& inst) { + return inst.GetOpcode() == opcode; + })); + Optimization::CollectShaderInfoPass(env, program); + program.info.native_global_bounds = true; + program.info.constant_buffer_descriptors = {{0, 1}}; + + Profile profile{}; + profile.supported_spirv = 0x00010300; + profile.support_int64 = true; + profile.support_descriptor_aliasing = true; + profile.unified_descriptor_binding = true; + profile.min_ssbo_alignment = 16; + return Backend::SPIRV::EmitSPIRV(profile, program); +} +} // namespace + +TEST_CASE("Untracked 32-bit global atomics route to their storage buffer", + "[shader_recompiler][global_memory]") { + constexpr u32 OpFunctionCall = 57; + constexpr u32 OpAtomicCompareExchange = 230; + // OpAtomicExchange 229, IAdd 234, SMin 236, UMin 237, SMax 238, UMax 239, And 240, Or 241, + // Xor 242. INC, DEC and F32 add go through the storage buffer CAS loops. + const std::array, 12> cases{{ + {IR::Opcode::GlobalAtomicIAdd32, 234}, + {IR::Opcode::GlobalAtomicSMin32, 236}, + {IR::Opcode::GlobalAtomicUMin32, 237}, + {IR::Opcode::GlobalAtomicSMax32, 238}, + {IR::Opcode::GlobalAtomicUMax32, 239}, + {IR::Opcode::GlobalAtomicAnd32, 240}, + {IR::Opcode::GlobalAtomicOr32, 241}, + {IR::Opcode::GlobalAtomicXor32, 242}, + {IR::Opcode::GlobalAtomicExchange32, 229}, + {IR::Opcode::GlobalAtomicInc32, OpAtomicCompareExchange}, + {IR::Opcode::GlobalAtomicDec32, OpAtomicCompareExchange}, + {IR::Opcode::GlobalAtomicAddF32, OpAtomicCompareExchange}, + }}; + for (const auto& [opcode, atomic] : cases) { + CAPTURE(static_cast(opcode)); + std::vector code; + REQUIRE_NOTHROW(code = EmitUntrackedAtomic(opcode)); + CHECK_FALSE(Instructions(code, atomic).empty()); + CHECK_FALSE(Instructions(code, OpFunctionCall).empty()); + } +} From a6dea8a75db8e080de793228d22e964b0cbc3798 Mon Sep 17 00:00:00 2001 From: zortos293 <65777760+zortos293@users.noreply.github.com> Date: Thu, 5 Nov 2026 00:38:18 +0000 Subject: [PATCH 08/11] shader: drop the U64 clamp from lowered SHF.R funnel immediates The Maxwell translator reads bits 20-38 of SHF.R imm as the shift, so the U64 max-shift bit made every lowered SHF.R.U64 shift by 63 (Ocarina vertex shader 1a55fa57a335b963 returned bit 63 instead of a 21-bit field). Ported from zortos293/Lichii fix/oot-bloom 4435506c96 without its test file (not in this tree); checked by a field test in ampere_lowering.cpp. In game it lowers Ocarina's gameplay brightness (mean 0.89 -> 0.82) but the bloom stays. Co-authored-by: zortos293 <65777760+zortos293@users.noreply.github.com> Assisted-by: Claude Opus 5.5 (ported the fix, wrote the test, ran the in-game check) (cherry picked from commit 223c752f442b2f8d11f703588fe5130c0faa8d10) --- .../frontend/ampere/lower_to_maxwell.h | 5 +++-- src/tests/video_core/ampere_lowering.cpp | 17 +++++++++++++++++ 2 files changed, 20 insertions(+), 2 deletions(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index c461c56b28..fe522294cd 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -2950,9 +2950,10 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 const bool is_signed = insn.Bits(73, 2) == 2; if (!insn.Bits(80, 1) && insn.Bits(76, 1) && insn.Bits(73, 2) == 1 && !insn.Bits(75, 1) && insn.words[1] <= 32) { + // SHF.R.U64 funnel: the translator reads bits 20-38 as the immediate, so the U64 + // clamp (bit 38) would turn every shift into 63; shifts up to 32 need no clamp. emit(base(0x38f8000000000000ULL) | (u64{insn.Bits(24, 8)} << 8) | - (u64{insn.words[1]} << 20) | (u64{2} << 37) | - (u64{insn.Bits(64, 8)} << 39)); + (u64{insn.words[1]} << 20) | (u64{insn.Bits(64, 8)} << 39)); break; } if (!insn.Bits(80, 1) && !insn.Bits(76, 1) && insn.Bits(64, 8) == 255 && diff --git a/src/tests/video_core/ampere_lowering.cpp b/src/tests/video_core/ampere_lowering.cpp index 0ea19ad9a6..21b9512d0b 100644 --- a/src/tests/video_core/ampere_lowering.cpp +++ b/src/tests/video_core/ampere_lowering.cpp @@ -102,6 +102,23 @@ TEST_CASE("Ampere SUST.P keeps its channel mask", "[video_core][ampere]") { REQUIRE(((FindMaxwell(lowered, 0xeb2, 52) >> 20) & 0xf) == 1); } +TEST_CASE("Ampere SHF.R.U64 immediate shifts keep their amount", "[video_core][ampere]") { + // SHF.R.U64 R14, R14, n, R17 (Ocarina vertex shader 1a55fa57a335b963). The Maxwell + // translator reads the immediate from bits 20-38, so no max-shift bits may be set there. + for (const u32 shift : {0U, 1U, 21U, 31U, 32U}) { + const std::array code{{ + {{0x0e0e7819, shift, 0x00001211, 0x020fe200}}, + EXIT, + }}; + const auto lowered = LowerToMaxwell(code); + REQUIRE(static_cast(lowered)); + const u64 shf = FindMaxwell(lowered, 0x38f8, 48); + REQUIRE(shf != 0); + CHECK(((shf >> 20) & 0x7ffff) == shift); + CHECK(((shf >> 39) & 0xff) == 17); + } +} + TEST_CASE("Ampere SHFL with a register lane keeps its mode", "[video_core][ampere]") { // SHFL.DOWN PT, R7, R0, R5, 0x1f const std::array code{{ From dc3e912fb76d0b318b4a14007d771017a3fb586b Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 04:26:51 +0100 Subject: [PATCH 09/11] shader: pack FP16 texture results with c0 in the low half Ampere returns a TEX.F16 result as packed pairs, c0 in the low half of the first destination. The lowering merges the two converted halves with ISCADD, which shifts op_a up by 16, but gave op_a the low half: H0 and H1 came out swapped. Ocarina of Time's tensor network reads blue from .H0, so its image lost blue (gameplay blue mean 0.03 -> 0.60 with EDEN_NX2_TENSOR=1). Ported from Quixomatic/Eglantine dev ee247a0b68 (Franseven). Co-authored-by: Franseven Assisted-by: Claude Opus 5.5 (found the fork commit, ported it, checked it in game) (cherry picked from commit 29ed33831a3191ea620211dcd03570992f7fdffb) --- src/shader_recompiler/frontend/ampere/lower_to_maxwell.h | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h index fe522294cd..55447c0240 100644 --- a/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/src/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -2439,7 +2439,9 @@ inline LoweredProgram LowerToMaxwellWith(std::span code, u32 if (target == 255) return; emit(0x5ca8000000000000ULL | predicate | 246 | (u64{1} << 8) | (u64{2} << 10) | (u64{low} << 20)); emit(0x5ca8000000000000ULL | predicate | 247 | (u64{1} << 8) | (u64{2} << 10) | (u64{high} << 20)); - emit(0x5c18000000000000ULL | predicate | target | (u64{246} << 8) | (u64{247} << 20) | (u64{16} << 39)); + // ISCADD shifts op_a (bits 8..15) up, so op_a carries the high half and + // op_b the low one; the other way round leaves H1 = c0 and H0 = c1. + emit(0x5c18000000000000ULL | predicate | target | (u64{247} << 8) | (u64{246} << 20) | (u64{16} << 39)); }; pack(static_cast(dest), 250, components > 1 ? 251U : 255U); if (components > 2) { From 8de2d42a2eefb67b97963de760b0041c20747629 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 04:26:51 +0100 Subject: [PATCH 10/11] buffer_cache: re-upload a native program's 64 KiB uniform bank once per frame Native programs keep tables next to their code in a 64 KiB bank read with indirect offsets. Writes to that memory can go unnoticed by the memory tracker, leaving stale host contents. Ported from Quixomatic/Eglantine dev 911309003d (Franseven, itself from monkeefunnee/Eglantine raiders-fix cbf7307b31), with the page constants qualified as in e732fd1290. Co-authored-by: Franseven Assisted-by: Claude Opus 5.5 (ported the fix) (cherry picked from commit 5c7da2713a85c4bdc78da004aa6a13577c87f14a) --- src/video_core/buffer_cache/buffer_cache.h | 18 ++++++++++++++++++ .../buffer_cache/buffer_cache_base.h | 3 +++ 2 files changed, 21 insertions(+) diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index f666520ac8..e4d6c2f628 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -106,6 +106,7 @@ void BufferCache

::TickFrame() { return; } runtime.TickFrame(slot_buffers); + native_indirect_uniforms_this_frame.clear(); if (frame_tick % 600 == 0 && !uploaded_buffer_ranges.empty()) { for (const auto& [address, size] : uploaded_buffer_ranges) { device_memory.HintCold(address, size); @@ -1138,6 +1139,23 @@ void BufferCache

::BindHostGraphicsUniformBuffer(size_t stage, u32 index, u32 return alignment > 1 && (offset % alignment) != 0; } }(); + // Native programs keep tables next to their code in a 64 KiB bank read with indirect offsets + // (Splatoon Raiders: the dialogue box corners in c1). Writes to that memory can go unnoticed by + // the memory tracker, leaving stale host contents, so upload such a bank again once per frame. + if (has_host_buffer && (*channel_state->uniform_buffer_sizes)[stage][index] == 0x10000 && + size != 0 && gpu_memory != nullptr && gpu_memory->UsesNativeTiling() && + !memory_tracker.IsRegionGpuModified(device_addr, size) && + native_indirect_uniforms_this_frame.insert(device_addr).second) { + // The bank size is only an upper limit: stop at the end of the mapped memory. + u64 mapped_size = 0; + while (mapped_size < size && + device_memory.template GetPointer(device_addr + mapped_size) != nullptr) { + mapped_size += Core::DEVICE_PAGESIZE - ((device_addr + mapped_size) & Core::DEVICE_PAGEMASK); + } + if (mapped_size != 0) { + memory_tracker.MarkRegionAsCpuModified(device_addr, (std::min)(mapped_size, size)); + } + } const bool use_fast_buffer = needs_alignment_stream || (has_host_buffer && size <= channel_state->uniform_buffer_skip_cache_size && !memory_tracker.IsRegionGpuModified(device_addr, size)); diff --git a/src/video_core/buffer_cache/buffer_cache_base.h b/src/video_core/buffer_cache/buffer_cache_base.h index ec442676d2..cc5a3aaf05 100644 --- a/src/video_core/buffer_cache/buffer_cache_base.h +++ b/src/video_core/buffer_cache/buffer_cache_base.h @@ -17,6 +17,7 @@ #include #include #include +#include #include #include "common/container/unordered_map.h" @@ -377,6 +378,8 @@ class BufferCache : public VideoCommon::ChannelSetupCaches native_indirect_uniforms_this_frame; private: template From b536a971a54524ef271be60baef830e368529797 Mon Sep 17 00:00:00 2001 From: kraibse <33869366+kraibse@users.noreply.github.com> Date: Thu, 5 Nov 2026 03:29:28 +0100 Subject: [PATCH 11/11] video_core: upscale stand-in takes its images from the native program When a tensor network is rejected, the stand-in blits its input to its output. Outside the known slots it took the first same-format pair in cbuf0+0x80..0x11c, which in Ocarina of Time's gameplay was a 4x4 image onto a 1920x1080 target and a 120x67 one over the 2560x1440 frame: the 3D scene went black. The handles the program itself reads (TEX/TLD/TLD4) and stores to (SUST.P) through cbuf0 are tried first, and every pair must keep the aspect ratio and scale by at most 4. Input and output may differ in format, and an E5B9G9R9 input (no 2D surface format) is blitted through a source format override in the texture cache, which on Vulkan samples the source through a plain view. Checked in game (hidden eden-cli, Kokiri Forest save): the scene draws again without EDEN_NX2_TENSOR, as bright as the network's own output. Assisted-by: Claude Opus 5.5 (found the wrong pairs, wrote the fix, ran the checks) (cherry picked from commit dbf3417099678469c68ebe411935d62df00f1677) --- src/video_core/engines/kepler_compute.cpp | 111 ++++++++++++++++-- src/video_core/engines/kepler_compute.h | 1 + src/video_core/texture_cache/texture_cache.h | 13 +- .../texture_cache/texture_cache_base.h | 4 + 4 files changed, 116 insertions(+), 13 deletions(-) diff --git a/src/video_core/engines/kepler_compute.cpp b/src/video_core/engines/kepler_compute.cpp index 23ffbad790..63d2b194eb 100644 --- a/src/video_core/engines/kepler_compute.cpp +++ b/src/video_core/engines/kepler_compute.cpp @@ -23,6 +23,7 @@ #include "video_core/nx2_profile.h" #include "video_core/nx2.h" #include "video_core/engines/fermi_2d.h" +#include "video_core/texture_cache/texture_cache_base.h" #include "video_core/engines/kepler_compute.h" #include "video_core/engines/maxwell_3d.h" #include "video_core/game_dlss_status.h" @@ -895,6 +896,44 @@ bool KeplerCompute::RememberDlssMotionVectors() { // 720p to 1440p docked). Blit one to the other with bilinear filtering when both are plain 2D // images of the same format and the output is larger: title 4's slots first, then Fortune's // Weave's, else the first two handles in cbuf0+0x80..0x11c that form such a pair. +// The cbuf0 byte offsets of the texture handles a native program reads (TEX, TLD, TLD4, TXD) and +// the surface handles it stores to (SUST.P), from its code up to the first unconditional EXIT. +std::pair, std::vector> KeplerCompute::NativeImageHandles(GPUVAddr program) { + static std::unordered_map, std::vector>> cache; + if (const auto it = cache.find(program); it != cache.end()) { + return it->second; + } + std::pair, std::vector> handles; + const auto add = [](std::vector& list, u32 offset) { + if (std::find(list.begin(), list.end(), offset) == list.end()) list.push_back(offset); + }; + std::array block{}; + bool exited = false; + for (u64 offset = 0; offset < 0x40000 && !exited; offset += sizeof(block)) { + if (!memory_manager.IsFullyMappedRange(program + offset, sizeof(block))) break; + memory_manager.ReadBlock(program + offset, block.data(), sizeof(block)); + for (size_t word = 0; word < block.size() && !exited; word += 4) { + const u32 opcode = block[word] & 0xfff; + const u32 binding = (block[word + 1] >> 22) & 0x1f; + const u32 handle = ((block[word + 1] >> 8) & 0x3fff) * 4; + switch (opcode) { + case 0xb60: case 0xb63: case 0xb66: case 0xb6d: + if (binding == 0) add(handles.first, handle); + break; + case 0xb9b: + if (binding == 0) add(handles.second, handle); + break; + case 0x94d: // EXIT, unconditional when guarded by PT + exited = ((block[word] >> 12) & 0xf) == 7; + break; + default: + break; + } + } + } + return cache.emplace(program, std::move(handles)).first->second; +} + bool KeplerCompute::UpscaleInsteadOfNativeProgram() { const auto& cbuf = launch_description.const_buffer_config[0]; const GPUVAddr cbuf_address = (GPUVAddr{cbuf.address_high} << 32) | cbuf.address_low; @@ -903,11 +942,19 @@ bool KeplerCompute::UpscaleInsteadOfNativeProgram() { if (tic_index > regs.tic.limit) return std::nullopt; return GetTICEntry(tic_index); }; + // An upscale keeps the aspect ratio and scales by at most 4: anything else (a 4x4 lookup + // table, a 120x67 luminance image) is another image the program happens to use. const auto upscale_pair = [](const std::optional& from, - const std::optional& to) { - return from && to && from->format == to->format && from->Address() != 0 && - to->Address() != 0 && to->Width() > from->Width() && from->IsBlockLinear() && - to->IsBlockLinear(); + const std::optional& to, bool same_format = true) { + if (!from || !to || (same_format && from->format != to->format) || from->Address() == 0 || + to->Address() == 0 || from->Address() == to->Address() || !from->IsBlockLinear() || + !to->IsBlockLinear()) + return false; + const u64 from_width = from->Width(), from_height = from->Height(); + const u64 to_width = to->Width(), to_height = to->Height(); + const u64 cross_a = to_width * from_height, cross_b = to_height * from_width; + return from_width >= 64 && to_width > from_width && to_width <= from_width * 4 && + (cross_a > cross_b ? cross_a - cross_b : cross_b - cross_a) * 32 <= cross_a; }; std::optional input = image(0x90), output = image(0xa0); bool fortunes_weave = false; @@ -917,6 +964,29 @@ bool KeplerCompute::UpscaleInsteadOfNativeProgram() { fortunes_weave = upscale_pair(input, output); } if (!upscale_pair(input, output)) { + // The images the program itself uses: textures it reads and the surface it stores to + // through cbuf0 handles (Ocarina of Time's gameplay network reads +0xc0, writes +0xd0). + // The input is the largest read image that the output upscales, the output's format first. + const auto [reads, stores] = NativeImageHandles(regs.code_loc.Address()); + input.reset(); + output.reset(); + for (const u32 store : stores) { + const auto to = image(store); + for (const bool same_format : {true, false}) { + for (const u32 read : reads) { + const auto from = image(read); + if (upscale_pair(from, to, same_format) && + (!input || u64{from->Width()} * from->Height() > u64{input->Width()} * input->Height())) { + input = from; + output = to; + } + } + if (input) break; + } + if (input) break; + } + } + if (!upscale_pair(input, output, false)) { input.reset(); output.reset(); for (u32 first = 0x80; first < 0x120 && !input; first += 4) { @@ -936,17 +1006,27 @@ bool KeplerCompute::UpscaleInsteadOfNativeProgram() { if (!input || !output) return false; using RTFormat = Tegra::RenderTargetFormat; - RTFormat format{}; - switch (input->format) { - case Texture::TextureFormat::A2B10G10R10: format = RTFormat::A2B10G10R10_UNORM; break; - case Texture::TextureFormat::A8B8G8R8: format = RTFormat::A8B8G8R8_UNORM; break; - case Texture::TextureFormat::B10G11R11: format = RTFormat::B10G11R11_FLOAT; break; - case Texture::TextureFormat::R16G16B16A16: format = RTFormat::R16G16B16A16_FLOAT; break; - default: return false; + const auto blit_format = [](const Texture::TICEntry& tic) -> std::optional { + switch (tic.format) { + case Texture::TextureFormat::A2B10G10R10: return RTFormat::A2B10G10R10_UNORM; + case Texture::TextureFormat::A8B8G8R8: + return tic.r_type == Texture::ComponentType::UNORM && tic.srgb_conversion + ? RTFormat::A8B8G8R8_SRGB + : RTFormat::A8B8G8R8_UNORM; + case Texture::TextureFormat::B10G11R11: return RTFormat::B10G11R11_FLOAT; + case Texture::TextureFormat::R16G16B16A16: return RTFormat::R16G16B16A16_FLOAT; + default: return std::nullopt; + } + }; + // E5B9G9R9 (Ocarina of Time's gameplay frame) has no 2D surface format: the blit reads it as a + // 32-bit placeholder surface and the texture cache samples it as E5B9G9R9. + const bool shared_exponent_input = input->format == Texture::TextureFormat::E5B9G9R9; + if ((!blit_format(*input) && !shared_exponent_input) || !blit_format(*output)) { + return false; } const auto surface = [&](const Texture::TICEntry& tic) { Fermi2D::Surface result{}; - result.format = format; + result.format = blit_format(tic).value_or(RTFormat::B10G11R11_FLOAT); result.linear = Fermi2D::MemoryLayout::BlockLinear; result.block_width.Assign(tic.block_width); result.block_height.Assign(tic.block_height); @@ -1042,6 +1122,9 @@ bool KeplerCompute::UpscaleInsteadOfNativeProgram() { } } auto source = surface(*input); + if (shared_exponent_input) { + source.format = RTFormat::B10G11R11_FLOAT; // same size per texel, replaced below + } // EDEN_NX2_UPSCALE_SOURCE=: blit from another same-sized image (debug aid). if (static const u64 override_source = [] { const char* value = std::getenv("EDEN_NX2_UPSCALE_SOURCE"); @@ -1071,8 +1154,12 @@ bool KeplerCompute::UpscaleInsteadOfNativeProgram() { dlss_depth ? &*dlss_depth : nullptr, dlss_jitter, {1.0f, 1.0f}); }; + if (shared_exponent_input) { + VideoCommon::g_blit_source_format = VideoCore::Surface::PixelFormat::E5B9G9R9_FLOAT; + } const bool done = upscale_temporally() || rasterizer->AccelerateSurfaceCopy(source, surface(*output), config); + VideoCommon::g_blit_source_format.reset(); if (static u32 checked{}; checked++ == sample_at) { const u64 size = u64{output->Width()} * output->Height() * 4; if (const auto device = memory_manager.GpuToCpuAddress(output->Address())) { diff --git a/src/video_core/engines/kepler_compute.h b/src/video_core/engines/kepler_compute.h index bda4898bc2..fa5c91adb4 100644 --- a/src/video_core/engines/kepler_compute.h +++ b/src/video_core/engines/kepler_compute.h @@ -278,6 +278,7 @@ class KeplerCompute final : public EngineInterface { void CheckDlssNetwork(); [[nodiscard]] VideoCore::GameDlssReason GameDlssBlocker() const; void UpdateGameDlssStatus(bool complete, VideoCore::GameDlssReason blocker); + std::pair, std::vector> NativeImageHandles(GPUVAddr program); bool UpscaleInsteadOfNativeProgram(); /// This frame's DLSS evaluation (Fortune's Weave): whether its passes all ran, and whether diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index 25a313dd22..3fe4629afe 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -1090,7 +1090,15 @@ bool TextureCache

::BlitImage(const Tegra::Engines::Fermi2D::Surface& dst, const std::optional src_base = src_image.TryFindBase(src.Address()); const SubresourceRange src_range{.base = src_base.value(), .extent = {1, 1}}; const ImageViewInfo src_view_info(ImageViewType::e2D, images.src_format, src_range); - const auto [src_framebuffer_id, src_view_id] = RenderTargetFromImage(src_id, src_view_info); + // Vulkan samples the source in a shader: a view is enough, and the source format need not be + // renderable (E5B9G9R9). + [[maybe_unused]] FramebufferId src_framebuffer_id; + ImageViewId src_view_id; + if constexpr (FRAMEBUFFER_BLITS) { + std::tie(src_framebuffer_id, src_view_id) = RenderTargetFromImage(src_id, src_view_info); + } else { + src_view_id = FindOrEmplaceImageView(src_id, src_view_info); + } const auto [src_samples_x, src_samples_y] = SamplesLog2(src_image.info.num_samples); Region2D src_region{ Offset2D{.x = copy.src_x0 >> src_samples_x, .y = copy.src_y0 >> src_samples_y}, @@ -2726,6 +2734,9 @@ std::optional::BlitImages> TextureCache

::GetBlitImag const GPUVAddr src_addr = src.Address(); ImageInfo dst_info(dst); ImageInfo src_info(src); + if (g_blit_source_format) { + src_info.format = *g_blit_source_format; + } const bool can_be_depth_blit = dst_info.format == src_info.format && copy.filter == Tegra::Engines::Fermi2D::Filter::Point; ImageId dst_id; diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h index 06fcc0de77..93deeca578 100644 --- a/src/video_core/texture_cache/texture_cache_base.h +++ b/src/video_core/texture_cache/texture_cache_base.h @@ -55,6 +55,10 @@ struct ChannelState; namespace VideoCommon { +/// Set around a 2D blit to read its source as this format instead of the surface's (for formats the +/// 2D engine has no surface format for, e.g. E5B9G9R9 by the NX2 upscale stand-in). GPU thread only. +inline thread_local std::optional g_blit_source_format; + using Tegra::Texture::TICEntry; using Tegra::Texture::TSCEntry; using VideoCore::Surface::PixelFormat;