monsterlag · Nov 8, 2026 · Nov 8, 2026 · Nov 9, 2026
diff --git a/.github/workflows/source-regressions.yml b/.github/workflows/source-regressions.yml
@@ -151,6 +151,35 @@ jobs:
             spirv-val --target-env vulkan1.3 --uniform-buffer-standard-layout "$module"
           done
 
+  spirv-f2i-immediate:
+    name: Native F2I immediate lowering
+    runs-on: ubuntu-24.04
+    timeout-minutes: 10
+    steps:
+      - name: Check out source
+        uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1
+        with:
+          persist-credentials: false
+      - name: Check out pinned SPIR-V headers
+        uses: actions/checkout@3d3c42e5aac5ba805825da76410c181273ba90b1 # v7.0.1
+        with:
+          repository: KhronosGroup/SPIRV-Headers
+          ref: 496543121ce6419f23d6fa5d7194ba66c36212d2
+          path: _deps/spirv-headers
+          persist-credentials: false
+      - name: Install Clang and validator
+        run: sudo apt-get update && sudo apt-get install -y clang spirv-tools
+      - name: Build the focused fixture
+        run: |
+          CC=clang CXX=clang++ cmake -S tests/f2i_immediate -B "${{ runner.temp }}/f2i" -DREPO_ROOT="${{ github.workspace }}" -DSPIRV_HEADERS_ROOT="${{ github.workspace }}/_deps/spirv-headers"
+          cmake --build "${{ runner.temp }}/f2i" --parallel 2
+      - name: Emit and validate authored modules
+        run: |
+          "${{ runner.temp }}/f2i/f2i_immediate" "${{ runner.temp }}/f2i/modules"
+          for module in "${{ runner.temp }}"/f2i/modules/*.spv; do
+            spirv-val --target-env vulkan1.3 --uniform-buffer-standard-layout "$module"
+          done
+
   actionlint:
     name: Workflow lint (actionlint)
     runs-on: ubuntu-24.04

diff --git a/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h b/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h
@@ -1959,7 +1959,8 @@ inline LoweredProgram LowerToMaxwellWith(std::span<const Instruction> code, u32
                  (u64{insn.Bits(32, 8)} << 20) | (u64{insn.Bits(78, 2)} << 39));
             break;
         }
-        case 0x305: case 0xb05: case 0xd05: { // F2I NTZ, register/constant-buffer/uniform source
+        case 0x305: case 0x905: case 0xb05: case 0xd05: { // F2I NTZ, register/immediate/constant-buffer/uniform source
+            const bool immediate = opcode == 0x905;
             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.
             const bool half = opcode == 0x305 && (mode & ~0x100U) == 0x0010f000;
@@ -1970,12 +1971,14 @@ inline LoweredProgram LowerToMaxwellWith(std::span<const Instruction> code, u32
             // (hookshot) a negated or absolute register source (bits 63, 62), made first
             const bool source_mod = opcode == 0x305 && !half && (insn.words[1] & 0x3fffff00U) == 0;
             if ((!half && !small && !floor && !nearest && !ceil && (mode & ~0x10100U) != 0x0020f000) ||
-                (opcode != 0xb05 && !source_mod && (insn.words[1] & (half ? 0xefffff00U : 0xffffff00U))) ||
+                (opcode != 0xb05 && !immediate && !source_mod && (insn.words[1] & (half ? 0xefffff00U : 0xffffff00U))) ||
+                (immediate && insn.Bits(24, 8)) ||
                 (opcode == 0xd05 && uniform_register(insn.Bits(32, 8)) == ~0U)) {
                 return fail(index, "native float-to-integer conversion mode");
             }
             const bool is_signed = insn.Bits(72, 1);
-            u32 source = opcode == 0xb05 ? materialize(2, 0, 252, false, false)
+            u32 source = immediate ? materialize(1, 0, 252, false, false)
+                       : opcode == 0xb05 ? materialize(2, 0, 252, false, false)
                        : opcode == 0xd05 ? uniform_register(insn.Bits(32, 8)) : insn.Bits(32, 8);
             if (source_mod && insn.Bits(62, 2)) source = materialize(0, insn.Bits(32, 8), 249, insn.Bits(62, 1), insn.Bits(63, 1));
             if (half) {

diff --git a/tests/f2i_immediate/CMakeLists.txt b/tests/f2i_immediate/CMakeLists.txt
@@ -0,0 +1,35 @@
+cmake_minimum_required(VERSION 3.20)
+project(f2i_immediate LANGUAGES CXX)
+set(CMAKE_CXX_STANDARD 20)
+set(CMAKE_CXX_STANDARD_REQUIRED ON)
+if(NOT DEFINED REPO_ROOT OR NOT DEFINED SPIRV_HEADERS_ROOT)
+    message(FATAL_ERROR "REPO_ROOT and SPIRV_HEADERS_ROOT are required")
+endif()
+file(TO_CMAKE_PATH "${REPO_ROOT}" REPO_ROOT)
+file(TO_CMAKE_PATH "${SPIRV_HEADERS_ROOT}" SPIRV_HEADERS_ROOT)
+set(SHADER "${REPO_ROOT}/gpu/vk/eden/shader_recompiler")
+file(GLOB_RECURSE BACKEND_SOURCES CONFIGURE_DEPENDS "${SHADER}/backend/spirv/*.cpp")
+file(GLOB_RECURSE IR_SOURCES CONFIGURE_DEPENDS "${SHADER}/frontend/ir/*.cpp")
+file(GLOB_RECURSE OPTIMIZATION_SOURCES CONFIGURE_DEPENDS "${SHADER}/ir_opt/*.cpp")
+file(GLOB_RECURSE SIRIT_SOURCES CONFIGURE_DEPENDS "${REPO_ROOT}/gpu/vk/sirit/src/*.cpp")
+set(IMPL "${SHADER}/frontend/maxwell/translate/impl")
+add_executable(f2i_immediate
+    "${REPO_ROOT}/tests/f2i_immediate/fixture.cpp"
+    "${SHADER}/frontend/maxwell/decode.cpp"
+    "${SHADER}/frontend/maxwell/opcodes.cpp" "${IMPL}/common_funcs.cpp"
+    "${IMPL}/impl.cpp" "${IMPL}/register_map_nx2.cpp" "${IMPL}/move_register.cpp"
+    "${IMPL}/floating_point_conversion_integer.cpp"
+    "${IMPL}/floating_point_compare_and_set.cpp" "${IMPL}/floating_point_min_max.cpp"
+    "${IMPL}/logic_operation.cpp"
+    ${BACKEND_SOURCES} ${IR_SOURCES} ${OPTIMIZATION_SOURCES} ${SIRIT_SOURCES})
+target_include_directories(f2i_immediate PRIVATE
+    "${REPO_ROOT}/gpu/vk/compat" "${REPO_ROOT}/gpu/compat"
+    "${REPO_ROOT}/gpu/vk/eden" "${REPO_ROOT}/gpu/third_party/suyu"
+    "${REPO_ROOT}/gpu/vk/sirit/include" "${REPO_ROOT}/gpu/vk/sirit/src"
+    "${SPIRV_HEADERS_ROOT}/include")
+target_compile_definitions(f2i_immediate PRIVATE NDEBUG MKWVK_INCLUDES=2 VK_NO_PROTOTYPES)
+if(MSVC)
+    target_compile_options(f2i_immediate PRIVATE /O1 /EHsc)
+else()
+    target_compile_options(f2i_immediate PRIVATE -O1)
+endif()
diff --git a/tests/f2i_immediate/README.md b/tests/f2i_immediate/README.md
@@ -0,0 +1,37 @@
+# Native F2I immediate regression
+
+This game-free fixture constructs its instructions from independently chosen
+registers, float values and bit patterns. It calls the production Ampere lowerer,
+Maxwell decoder and conversion implementations, then emits real IR and SPIR-V.
+No game, dump, cache, driver or GPU is used.
+
+The immediate contains all 32 source bits. Bits 62 and 63 of the native
+instruction belong to that literal; they are not the register source's abs/neg
+modifiers. The test checks literal preservation, signedness, destination width,
+four rounding modes, FTZ, ordinary/inverted predicates and `@!PT` skipping. It
+also checks existing register/cbuf/uniform source forms and explicitly rejects
+immediate F16, F64, 64-bit output, non-NTZ and unknown modes. The accepted U16
+form remains FTZ/nearest only.
+
+```sh
+CC=clang CXX=clang++ cmake -S tests/f2i_immediate -B build-f2i \
+  -DREPO_ROOT="$PWD" -DSPIRV_HEADERS_ROOT=/path/to/SPIRV-Headers
+cmake --build build-f2i --parallel 2
+build-f2i/f2i_immediate build-f2i/modules
+for module in build-f2i/modules/*.spv; do
+  spirv-val --target-env vulkan1.3 --uniform-buffer-standard-layout "$module"
+done
+```
+
+The unmodified `232eecb` lowerer fails the first authored immediate case with
+`unsupported native opcode`. The fixture tests the conversion path and its IR
+controls; SPIR-V validation does not establish numerical results on a GPU or a
+fix for any particular visible game symptom. In particular, it does not change
+or independently certify existing register-path saturation at numeric limits.
+
+Encoding reference: [Mesa NAK, pinned sm70_encode.rs](https://chromium.googlesource.com/external/gitlab.freedesktop.org/mesa/mesa/+/b1fe47e94409d81924c8663f58fac97e66ab4527/src/nouveau/compiler/nak/sm70_encode.rs#2134)
+(`OpF2I`, `encode_alu_base` and `encode_alu_imm`). The `0x105` base plus immediate
+form 4 in bits 9–11 gives `0x905`. Inputs or outputs wider than 32 bits use a
+different base, `0x111`. [NVIDIA PTX conversion documentation](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvt)
+describes rounding, FTZ, saturation and NaN results; it is not a SASS encoding
+specification.
diff --git a/tests/f2i_immediate/fixture.cpp b/tests/f2i_immediate/fixture.cpp
@@ -0,0 +1,227 @@
+// Authored instructions only. No dump, native program, GPU or game is read.
+#include <array>
+#include <bit>
+#include <cstdio>
+#include <cstdlib>
+#include <filesystem>
+#include <fstream>
+#include <optional>
+#include <string>
+#include <vector>
+#include <spirv/unified1/spirv.hpp>
+#include "shader_recompiler/backend/spirv/emit_spirv.h"
+#include "shader_recompiler/environment.h"
+#include "shader_recompiler/frontend/ampere/lower_to_maxwell.h"
+#include "shader_recompiler/frontend/maxwell/decode.h"
+#include "shader_recompiler/frontend/maxwell/translate/impl/impl.h"
+#include "shader_recompiler/ir_opt/passes.h"
+#include "shader_recompiler/object_pool.h"
+
+using namespace Shader;
+using Native = Ampere::Instruction;
+#define CHECK(x) do { if (!(x)) { std::fprintf(stderr, "FAIL:%d %s\n", __LINE__, #x); std::exit(1); } } while (0)
+
+// Mesa NAK sm70_encode.rs: OpF2I base 0x105, immediate form 4,
+// source bits 32..63; sign 72, destination width 75..76, NTZ 77,
+// rounding 78..79, FTZ 80, source width 84..85.
+static u32 mode(bool sign, u32 rounding, bool ftz, u32 dst = 2, u32 src = 2) {
+    return (u32{sign} << 8) | (dst << 11) | (1U << 13) |
+           (rounding << 14) | (u32{ftz} << 16) | (src << 20);
+}
+static Native instruction(u32 op, u32 dest, u32 source, u32 flags, u32 pred = 7) {
+    return {{{op | (pred << 12) | (dest << 16), source, flags, 0}}};
+}
+static Ampere::LoweredProgram lower(Native i) {
+    const std::array code{i, instruction(0x94d, 0, 0, 0x03800000)};
+    return Ampere::LowerToMaxwell(code);
+}
+
+struct FixtureEnvironment final : Shader::Environment {
+    FixtureEnvironment() { stage = Stage::Compute; }
+    u64 ReadInstruction(u32) override { std::abort(); }
+    u32 ReadCbufValue(u32, u32) override { return 0; }
+    TextureType ReadTextureType(u32) override { return TextureType::Color2D; }
+    TexturePixelFormat ReadTexturePixelFormat(u32) override { return TexturePixelFormat::R8_UNORM; }
+    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<u32, 3> WorkgroupSize() const override { return {1, 1, 1}; }
+    bool HasHLEMacroState() const override { return false; }
+    std::optional<ReplaceConstant> GetReplaceConstBuffer(u32, u32) override { return std::nullopt; }
+    void Dump(u64, u64) override { std::abort(); }
+};
+
+// Dispatch the real decoder's results to the actual implementations used by
+// this lowering. This fixture does not reimplement conversion arithmetic.
+static void translate(Maxwell::TranslatorVisitor& visitor, u64 word) {
+    switch (Maxwell::Decode(word)) {
+    case Maxwell::Opcode::MOV32I: visitor.MOV32I(word); break;
+    case Maxwell::Opcode::F2I_reg: visitor.F2I_reg(word); break;
+    case Maxwell::Opcode::FSET_reg: visitor.FSET_reg(word); break;
+    case Maxwell::Opcode::FMNMX_reg: visitor.FMNMX_reg(word); break;
+    case Maxwell::Opcode::LOP_reg: visitor.LOP_reg(word); break;
+    case Maxwell::Opcode::EXIT: break;
+    default: CHECK(false);
+    }
+}
+
+static unsigned modules = 0;
+static void check_pipeline(const std::filesystem::path& output, u32 literal, bool sign,
+                           u32 rounding, bool ftz, u32 dst_width = 2) {
+    constexpr u32 dest = 13;
+    const auto lowered = lower(instruction(0x905, dest, literal, mode(sign, rounding, ftz, dst_width)));
+    if (!lowered) std::fprintf(stderr, "native F2I immediate: %.*s at +%u\n",
+                              int(lowered.error.size()), lowered.error.data(), lowered.error_offset);
+    CHECK(lowered);
+    CHECK(lowered.instructions.size() > 2);
+    const u64 first = lowered.instructions.front();
+    CHECK(Maxwell::Decode(first) == Maxwell::Opcode::MOV32I);
+    CHECK(u32(first >> 20) == literal); // all 32 bits, including former modifier positions
+    const u32 scratch = first & 255;
+    CHECK(scratch >= 246 && scratch < 255);
+    bool converted = false;
+    for (u64 word : lowered.instructions) {
+        CHECK(((word >> 16) & 15) == 7);
+        CHECK(Maxwell::Decode(word) != Maxwell::Opcode::F2I_imm);
+        if (Maxwell::Decode(word) == Maxwell::Opcode::F2I_reg) {
+            CHECK((word & 255) == dest);
+            CHECK(((word >> 8) & 3) == dst_width);
+            CHECK(((word >> 10) & 3) == 2);
+            CHECK(((word >> 12) & 1) == u32(sign));
+            CHECK(((word >> 39) & 3) == rounding);
+            CHECK(((word >> 44) & 1) == u32(ftz));
+            CHECK(((word >> 45) & 1) == 0 && ((word >> 49) & 1) == 0);
+            if (sign) CHECK(((word >> 20) & 255) == scratch);
+            converted = true;
+        }
+    }
+    CHECK(converted);
+
+    ObjectPool<IR::Inst> pool;
+    IR::Block block{pool};
+    IR::Program program;
+    program.stage = Stage::Compute;
+    program.workgroup_size = {1, 1, 1};
+    program.shared_memory_size = 4;
+    program.blocks.push_back(&block);
+    program.post_order_blocks.push_back(&block);
+    IR::AbstractSyntaxNode node{};
+    node.type = IR::AbstractSyntaxNode::Type::Block;
+    node.data.block = &block;
+    program.syntax_list.push_back(node);
+    IR::AbstractSyntaxNode ret{};
+    ret.type = IR::AbstractSyntaxNode::Type::Return;
+    program.syntax_list.push_back(ret);
+    FixtureEnvironment env;
+    Maxwell::TranslatorVisitor visitor{env, block};
+    for (u64 word : lowered.instructions) translate(visitor, word);
+    visitor.ir.WriteShared(32, visitor.ir.Imm32(0), visitor.X(static_cast<IR::Reg>(dest)));
+    visitor.ir.Epilogue();
+    const IR::Opcode expected_round[] = {IR::Opcode::FPRoundEven32, IR::Opcode::FPFloor32,
+                                         IR::Opcode::FPCeil32, IR::Opcode::FPTrunc32};
+    bool saw_round = false, saw_convert = false, saw_nan_zero = false;
+    for (const IR::Inst& inst : block.Instructions()) {
+        if (inst.GetOpcode() == expected_round[rounding]) {
+            CHECK(inst.Flags<IR::FpControl>().fmz_mode == (ftz ? IR::FmzMode::FTZ : IR::FmzMode::None));
+            saw_round = true;
+        }
+        saw_convert |= inst.GetOpcode() == (sign ? IR::Opcode::ConvertS32F32 : IR::Opcode::ConvertU32F32);
+        if (inst.GetOpcode() == IR::Opcode::SelectU32 && !inst.Arg(0).IsImmediate() &&
+            inst.Arg(0).InstRecursive()->GetOpcode() == IR::Opcode::FPIsNan32 &&
+            inst.Arg(1).IsImmediate() && inst.Arg(1).U32() == 0)
+            saw_nan_zero = true;
+    }
+    CHECK(saw_round && saw_convert);
+    if (sign) CHECK(saw_nan_zero); // existing NTZ conversion explicitly selects zero for NaN
+    Optimization::SsaRewritePass(program);
+    Optimization::ConstantPropagationPass(env, program);
+    Optimization::DeadCodeEliminationPass(program);
+    Optimization::IdentityRemovalPass(program);
+    Optimization::CollectShaderInfoPass(env, program);
+    Profile profile{};
+    profile.supported_spirv = 0x00010600;
+    const auto spv = Backend::SPIRV::EmitSPIRV(profile, program);
+    CHECK(spv.size() > 5);
+    bool spirv_conversion = false, spirv_store = false;
+    for (std::size_t at = 5; at < spv.size();) {
+        const u32 count = spv[at] >> 16;
+        CHECK(count && at + count <= spv.size());
+        const auto op = static_cast<spv::Op>(spv[at] & 0xffff);
+        spirv_conversion |= op == (sign ? spv::Op::OpConvertFToS : spv::Op::OpConvertFToU);
+        spirv_store |= op == spv::Op::OpStore;
+        at += count;
+    }
+    CHECK(spirv_conversion && spirv_store); // result survives the production cleanup passes
+    const auto path = output / ("case-" + std::to_string(modules++) + ".spv");
+    std::ofstream stream(path, std::ios::binary);
+    stream.write(reinterpret_cast<const char*>(spv.data()), spv.size() * sizeof(u32));
+    CHECK(stream.good());
+}
+
+static void check_source_controls() {
+    // Existing source forms remain accepted with their original source handling.
+    for (u32 op : {0x305U, 0xb05U, 0xd05U}) {
+        const u32 source = op == 0xb05 ? 0x100U << 6 : 19U;
+        const u32 flags = mode(true, 3, false) | (op == 0xd05 ? 1U << 27 : 0);
+        const auto result = lower(instruction(op, 13, source, flags));
+        CHECK(result);
+        bool converted = false;
+        for (u64 word : result.instructions) converted |= Maxwell::Decode(word) == Maxwell::Opcode::F2I_reg;
+        CHECK(converted);
+        u64 hash = 14695981039346656037ULL;
+        for (u64 word : result.instructions) hash = (hash ^ word) * 1099511628211ULL;
+        std::printf("control %03x: %016llx\n", op, static_cast<unsigned long long>(hash));
+    }
+    // Existing register F16 support is preserved; adding an immediate does not
+    // silently add half-precision, double-precision, 64-bit output or non-NTZ modes.
+    CHECK(lower(instruction(0x305, 13, 19, mode(true, 3, false, 2, 1))));
+}
+
+static void check_controls() {
+    check_source_controls();
+    CHECK(!lower(instruction(0x905, 13, 0x3f800000, mode(true, 3, false, 2, 1))));
+    CHECK(!lower(instruction(0x905, 13, 0x3f800000, mode(true, 3, false, 2, 3))));
+    CHECK(!lower(instruction(0x905, 12, 0x3f800000, mode(false, 3, true, 3))));
+    CHECK(!lower(instruction(0x911, 12, 0x3f800000, mode(false, 3, true, 3))));
+    CHECK(!lower(instruction(0x905, 13, 0x3f800000, mode(true, 3, false) & ~(1U << 13))));
+    CHECK(!lower(instruction(0x905, 13, 0x3f800000, mode(true, 3, false) | 1U)));
+    auto unused_source = instruction(0x905, 13, 0x3f800000, mode(true, 3, false));
+    unused_source.words[0] |= 1U << 24;
+    CHECK(!lower(unused_source));
+    // The materialization and conversion keep ordinary and inverted guards.
+    for (bool sign : {false, true})
+        for (u32 pred : {0U, 10U, 7U}) {
+            const auto result = lower(instruction(0x905, 13, 0xc1540001, mode(sign, 3, false), pred));
+            CHECK(result);
+            for (std::size_t n = 0; n + 1 < result.instructions.size(); ++n)
+                CHECK(((result.instructions[n] >> 16) & 15) == pred);
+        }
+    const auto skipped = lower(instruction(0x905, 13, 0xffffffff, 0xffffffff, 15));
+    CHECK(skipped && skipped.instructions.size() == 1);
+    CHECK(Maxwell::Decode(skipped.instructions.front()) == Maxwell::Opcode::EXIT);
+}
+
+int main(int argc, char** argv) {
+    CHECK(argc == 2);
+    if (std::string(argv[1]) == "--existing-controls") {
+        check_source_controls();
+        return 0;
+    }
+    const std::filesystem::path output{argv[1]};
+    std::filesystem::create_directories(output);
+    // Bit-casts from authored finite values plus chosen zeros, subnormals,
+    // infinities and NaN payloads; none originate in a program dump.
+    const std::array literals{std::bit_cast<u32>(13.75f), std::bit_cast<u32>(-9.25f),
+        0U, 0x80000000U, 1U, 0x80000001U, 0x7f800000U, 0xff800000U, 0x7fc12345U, 0xffa54321U};
+    for (bool sign : {false, true})
+        for (u32 rounding = 0; rounding < 4; ++rounding)
+            for (bool ftz : {false, true})
+                check_pipeline(output, literals[(u32{sign} * 8 + rounding * 2 + u32{ftz}) % literals.size()],
+                               sign, rounding, ftz);
+    for (u32 literal : literals) check_pipeline(output, literal, true, 3, true);
+    check_pipeline(output, std::bit_cast<u32>(32768.5f), false, 0, true, 1);
+    check_controls();
+    std::printf("PASS: %u authored native F2I immediate -> Maxwell -> IR -> SPIR-V cases; source/mode/predicate controls pass\n", modules);
+}