From dcd471163e023c1cbd7da72aacf7ce70513eaacf Mon Sep 17 00:00:00 2001 From: monsterlag <64719697+monsterlag@users.noreply.github.com> Date: Sun, 8 Nov 2026 17:37:00 -0500 Subject: [PATCH] Translate native F2I instructions with immediate operands --- .github/workflows/source-regressions.yml | 29 +++ .../frontend/ampere/lower_to_maxwell.h | 9 +- tests/f2i_immediate/CMakeLists.txt | 35 +++ tests/f2i_immediate/README.md | 37 +++ tests/f2i_immediate/fixture.cpp | 227 ++++++++++++++++++ 5 files changed, 334 insertions(+), 3 deletions(-) create mode 100644 tests/f2i_immediate/CMakeLists.txt create mode 100644 tests/f2i_immediate/README.md create mode 100644 tests/f2i_immediate/fixture.cpp diff --git a/.github/workflows/source-regressions.yml b/.github/workflows/source-regressions.yml index 42bdf63f0..e1cb2fde3 100644 --- a/.github/workflows/source-regressions.yml +++ b/.github/workflows/source-regressions.yml @@ -147,6 +147,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 index 13e80bcf4..c2d877660 100644 --- a/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h +++ b/gpu/vk/eden/shader_recompiler/frontend/ampere/lower_to_maxwell.h @@ -1953,7 +1953,8 @@ inline LoweredProgram LowerToMaxwellWith(std::span 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; @@ -1964,12 +1965,14 @@ inline LoweredProgram LowerToMaxwellWith(std::span 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 new file mode 100644 index 000000000..47e5e717d --- /dev/null +++ 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 new file mode 100644 index 000000000..919dfd669 --- /dev/null +++ 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 new file mode 100644 index 000000000..dae9d6426 --- /dev/null +++ b/tests/f2i_immediate/fixture.cpp @@ -0,0 +1,227 @@ +// Authored instructions only. No dump, native program, GPU or game is read. +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#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 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 { 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 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 = █ + 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(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().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[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(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(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(13.75f), std::bit_cast(-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(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); +}