tool for ir dumps

CMakeLists.txt

@@ -784,6 +784,7 @@ add_subdirectory(src)
 if (ENABLE_DEBUG_TOOLS)
     add_subdirectory(tools/maxwell-disas)
     add_subdirectory(tools/maxwell-spirv)
+    add_subdirectory(tools/maxwell-ir)
 endif()
 
 # Set yuzu project or yuzu-cmd project as default StartUp Project in Visual Studio depending on whether QT is enabled or not

tools/README.md

@@ -9,7 +9,8 @@ Tools for Eden and other subprojects.
 ## Binaries
 
 - `maxwell-spirv`: Converts Maxwell shaders (dumped from `.ash` files) into SPIR-V code (emitted into STDOUT).
-- `maxwell-disas`: Dumb maxwell dissasembler.
+- `maxwell-disas`: Dumb raw Maxwell dissasembler.
+- `maxwell-ir`: Dump generated IR of Maxwell shaders.
 
 ## Scripts
 

tools/maxwell-disas/CMakeLists.txt

@@ -1,6 +1,8 @@
 # SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
 # SPDX-License-Identifier: GPL-3.0-or-later
-add_executable(maxwell-disas main.cpp)
+add_executable(maxwell-disas
+    main.cpp
+)
 target_link_libraries(maxwell-disas PRIVATE common shader_recompiler Threads::Threads)
 target_include_directories(maxwell-disas PRIVATE ${CMAKE_SOURCE_DIR}/src)
 if(UNIX AND NOT APPLE)

tools/maxwell-disas/file_environment.h

@@ -0,0 +1,60 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#pragma once
+
+#include <cstdlib>
+#include <sys/stat.h>
+#include "shader_recompiler/environment.h"
+
+class FileEnvironment final : public Shader::Environment {
+public:
+    FileEnvironment() = default;
+    ~FileEnvironment() override = default;
+    FileEnvironment& operator=(FileEnvironment&&) noexcept = default;
+    FileEnvironment(FileEnvironment&&) noexcept = default;
+    FileEnvironment& operator=(const FileEnvironment&) = delete;
+    FileEnvironment(const FileEnvironment&) = delete;
+    void Deserialize(std::ifstream& file) {}
+    [[nodiscard]] u64 ReadInstruction(u32 address) override {
+        if (address < read_lowest || address > read_highest) {
+            std::printf("cant read %08x\n", address);
+            std::abort();
+        }
+        return code[(address - read_lowest) / sizeof(u64)];
+    }
+    [[nodiscard]] u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override { return 0; }
+    [[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override {
+        auto const it{texture_types.find(handle)};
+        return it->second;
+    }
+    [[nodiscard]] Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override {
+        auto const it{texture_pixel_formats.find(handle)};
+        return it->second;
+    }
+    [[nodiscard]] bool IsTexturePixelFormatInteger(u32 handle) override { return true; }
+    [[nodiscard]] u32 ReadViewportTransformState() override { return viewport_transform_state; }
+    [[nodiscard]] u32 LocalMemorySize() const override { return local_memory_size; }
+    [[nodiscard]] u32 SharedMemorySize() const override { return shared_memory_size; }
+    [[nodiscard]] u32 TextureBoundBuffer() const override { return texture_bound; }
+    [[nodiscard]] std::array<u32, 3> WorkgroupSize() const override { return workgroup_size; }
+    [[nodiscard]] std::optional<Shader::ReplaceConstant> GetReplaceConstBuffer(u32 bank, u32 offset) override {
+        auto const it = cbuf_replacements.find((u64(bank) << 32) | u64(offset));
+        return it != cbuf_replacements.end() ? std::optional{it->second} : std::nullopt;
+    }
+    [[nodiscard]] bool HasHLEMacroState() const override { return cbuf_replacements.size() != 0; }
+    void Dump(u64 pipeline_hash, u64 shader_hash) override {}
+    std::vector<u64> code;
+    std::unordered_map<u32, Shader::TextureType> texture_types;
+    std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
+    std::unordered_map<u64, u32> cbuf_values;
+    std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;
+    std::array<u32, 3> workgroup_size{};
+    u32 local_memory_size{};
+    u32 shared_memory_size{};
+    u32 texture_bound{};
+    u32 read_lowest{};
+    u32 read_highest{};
+    u32 initial_offset{};
+    u32 viewport_transform_state = 1;
+};

tools/maxwell-disas/main.cpp

@@ -162,7 +162,7 @@ std::string_view SpecialRegGetName(size_t i) {
 }
 #include "generated.cpp"
 
-int ReferenceImpl(int argc, char *argv[]) {
+int DisasReferenceImpl(int argc, char *argv[]) {
     std::vector<uint64_t> code;
     FILE *fp = fopen(argv[1], "rb");
     if (fp != NULL) {
@@ -181,7 +181,7 @@ int ReferenceImpl(int argc, char *argv[]) {
     return EXIT_SUCCESS;
 }
 
-int ShaderRecompilerDisas(int argc, char *argv[]) {
+int DisasShaderRecompilerImpl(int argc, char *argv[]) {
     std::vector<u64> code;
     FILE *fp = fopen(argv[1], "rb");
     if (fp != NULL) {
@@ -203,17 +203,19 @@ int ShaderRecompilerDisas(int argc, char *argv[]) {
 int main(int argc, char *argv[]) {
     if (argc < 2) {
         printf(
-            "usage: %s [input file] [-n]\n"
+            "usage: %s [input file] [-n/i]\n"
             "Specify -n to use a disassembler that is NOT tied to the shader recompiler\n"
             "aka. a reference disassembler\n"
             , argv[0]);
         return EXIT_FAILURE;
     }
+
+    //DumpProgram
+
     if (argc >= 3) {
-        if (::strcmp(argv[2], "-n") == 0
+        if (::strcmp(argv[2], "-n") == 0 || ::strcmp(argv[2], "--new") == 0) {
-        || ::strcmp(argv[2], "--new") == 0) {
+            return DisasReferenceImpl(argc, argv);
-            return ReferenceImpl(argc, argv);
         }
     }
-    return ShaderRecompilerDisas(argc, argv);
+    return DisasShaderRecompilerImpl(argc, argv);
 }

tools/maxwell-ir/CMakeLists.txt

@@ -0,0 +1,9 @@
+# SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+# SPDX-License-Identifier: GPL-3.0-or-later
+add_executable(maxwell-ir main.cpp)
+target_link_libraries(maxwell-ir PRIVATE common shader_recompiler Threads::Threads)
+target_include_directories(maxwell-ir PRIVATE ${CMAKE_SOURCE_DIR}/src)
+if(UNIX AND NOT APPLE)
+    install(TARGETS maxwell-ir)
+endif()
+create_target_directory_groups(maxwell-ir)

tools/maxwell-ir/main.cpp

@@ -0,0 +1,59 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+
+#include "shader_recompiler/frontend/ir/basic_block.h"
+#include "shader_recompiler/frontend/maxwell/control_flow.h"
+#include "shader_recompiler/frontend/maxwell/translate_program.h"
+#include "shader_recompiler/host_translate_info.h"
+#include "shader_recompiler/object_pool.h"
+#include "../maxwell-disas/file_environment.h"
+
+int IrShaderRecompilerImpl(int argc, char *argv[]) {
+    size_t cfg_offset = 0;
+
+    Shader::ObjectPool<Shader::IR::Inst> inst_pool;
+    Shader::ObjectPool<Shader::IR::Block> block_pool;
+    Shader::ObjectPool<Shader::Maxwell::Flow::Block> cfg_blocks;
+    FileEnvironment env;
+
+    FILE *fp = fopen(argv[1], "rb");
+    if (fp != NULL) {
+        struct stat st;
+        fstat(fileno(fp), &st);
+        auto const words = (st.st_size / sizeof(u64));
+        env.code.resize(words + 1);
+        fread(env.code.data(), sizeof(u64), words, fp);
+        fclose(fp);
+    }
+
+    env.read_highest = env.read_lowest + env.code.size() * sizeof(u64);
+
+    Shader::Maxwell::Flow::CFG cfg(env, cfg_blocks, cfg_offset);
+
+    Shader::HostTranslateInfo host_info;
+    host_info.support_float64 = true;
+    host_info.support_float16 = true;
+    host_info.support_int64 = true;
+    host_info.needs_demote_reorder = true;
+    host_info.support_snorm_render_buffer = true;
+    host_info.support_viewport_index_layer = true;
+    host_info.support_geometry_shader_passthrough = true;
+    host_info.support_conditional_barrier = true;
+    host_info.min_ssbo_alignment = 0;
+    auto program = Shader::Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg, host_info);
+    auto const dumped_ir = Shader::IR::DumpProgram(program);
+    std::printf("%s\n", dumped_ir.c_str());
+    return EXIT_SUCCESS;
+}
+
+int main(int argc, char *argv[]) {
+    if (argc < 2) {
+        printf("usage: %s [input file]\n", argv[0]);
+        return EXIT_FAILURE;
+    }
+    return IrShaderRecompilerImpl(argc, argv);
+}

tools/maxwell-spirv/CMakeLists.txt

@@ -2,8 +2,8 @@
 # SPDX-License-Identifier: GPL-3.0-or-later
 add_executable(maxwell-spirv
     main.cpp
-    recompiler_impl.cpp
+    spirv_recompiler_impl.cpp
-    reference_impl.cpp
+    spirv_reference_impl.cpp
 )
 target_link_libraries(maxwell-spirv PRIVATE common shader_recompiler Threads::Threads)
 target_include_directories(maxwell-spirv PRIVATE ${CMAKE_SOURCE_DIR}/src)

tools/maxwell-spirv/main.cpp

@@ -5,8 +5,8 @@
 #include <cstdlib>
 #include <cstring>
 
-int ReferenceImpl(int argc, char *argv[]);
+int SpirvReferenceImpl(int argc, char *argv[]);
-int ShaderRecompilerImpl(int argc, char *argv[]);
+int SpirvShaderRecompilerImpl(int argc, char *argv[]);
 
 int main(int argc, char *argv[]) {
     if (argc < 2) {
@@ -19,8 +19,8 @@ int main(int argc, char *argv[]) {
     if (argc >= 3) {
         if (::strcmp(argv[2], "-n") == 0
         || ::strcmp(argv[2], "--new") == 0) {
-            return ReferenceImpl(argc, argv);
+            return SpirvReferenceImpl(argc, argv);
         }
     }
-    return ShaderRecompilerImpl(argc, argv);
+    return SpirvShaderRecompilerImpl(argc, argv);
 }

tools/maxwell-spirv/recompiler_impl.cpp

@@ -1,162 +0,0 @@
-// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
-// SPDX-License-Identifier: GPL-3.0-or-later
-
-#include <cstdlib>
-#include <sys/stat.h>
-#include "shader_recompiler/backend/spirv/emit_spirv.h"
-#include "shader_recompiler/environment.h"
-#include "shader_recompiler/frontend/maxwell/control_flow.h"
-#include "shader_recompiler/frontend/maxwell/translate_program.h"
-#include "shader_recompiler/host_translate_info.h"
-#include "shader_recompiler/object_pool.h"
-#include "shader_recompiler/profile.h"
-#include "shader_recompiler/runtime_info.h"
-
-class FileEnvironment final : public Shader::Environment {
-public:
-    FileEnvironment() = default;
-    ~FileEnvironment() override = default;
-    FileEnvironment& operator=(FileEnvironment&&) noexcept = default;
-    FileEnvironment(FileEnvironment&&) noexcept = default;
-    FileEnvironment& operator=(const FileEnvironment&) = delete;
-    FileEnvironment(const FileEnvironment&) = delete;
-    void Deserialize(std::ifstream& file);
-    [[nodiscard]] u64 ReadInstruction(u32 address) override;
-    [[nodiscard]] u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override;
-    [[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override;
-    [[nodiscard]] Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
-    [[nodiscard]] bool IsTexturePixelFormatInteger(u32 handle) override;
-    [[nodiscard]] u32 ReadViewportTransformState() override;
-    [[nodiscard]] u32 LocalMemorySize() const override;
-    [[nodiscard]] u32 SharedMemorySize() const override;
-    [[nodiscard]] u32 TextureBoundBuffer() const override;
-    [[nodiscard]] std::array<u32, 3> WorkgroupSize() const override;
-    [[nodiscard]] std::optional<Shader::ReplaceConstant> GetReplaceConstBuffer(u32 bank, u32 offset) override;
-    [[nodiscard]] bool HasHLEMacroState() const override {
-        return cbuf_replacements.size() != 0;
-    }
-    void Dump(u64 pipeline_hash, u64 shader_hash) override;
-
-    std::vector<u64> code;
-    std::unordered_map<u32, Shader::TextureType> texture_types;
-    std::unordered_map<u32, Shader::TexturePixelFormat> texture_pixel_formats;
-    std::unordered_map<u64, u32> cbuf_values;
-    std::unordered_map<u64, Shader::ReplaceConstant> cbuf_replacements;
-    std::array<u32, 3> workgroup_size{};
-    u32 local_memory_size{};
-    u32 shared_memory_size{};
-    u32 texture_bound{};
-    u32 read_lowest{};
-    u32 read_highest{};
-    u32 initial_offset{};
-    u32 viewport_transform_state = 1;
-};
-
-void FileEnvironment::Deserialize(std::ifstream& file) {}
-
-void FileEnvironment::Dump(u64 pipeline_hash, u64 shader_hash) {
-    //DumpImpl(pipeline_hash, shader_hash, code, read_highest, read_lowest, initial_offset, stage);
-}
-
-u64 FileEnvironment::ReadInstruction(u32 address) {
-    if (address < read_lowest || address > read_highest) {
-        std::printf("cant read %08x\n", address);
-        std::abort();
-    }
-    return code[(address - read_lowest) / sizeof(u64)];
-}
-
-u32 FileEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
-    return 0;
-}
-
-Shader::TextureType FileEnvironment::ReadTextureType(u32 handle) {
-    auto const it{texture_types.find(handle)};
-    return it->second;
-}
-
-Shader::TexturePixelFormat FileEnvironment::ReadTexturePixelFormat(u32 handle) {
-    auto const it{texture_pixel_formats.find(handle)};
-    return it->second;
-}
-
-bool FileEnvironment::IsTexturePixelFormatInteger(u32 handle) {
-    return true;
-}
-
-u32 FileEnvironment::ReadViewportTransformState() {
-    return viewport_transform_state;
-}
-
-u32 FileEnvironment::LocalMemorySize() const {
-    return local_memory_size;
-}
-
-u32 FileEnvironment::SharedMemorySize() const {
-    return shared_memory_size;
-}
-
-u32 FileEnvironment::TextureBoundBuffer() const {
-    return texture_bound;
-}
-
-std::array<u32, 3> FileEnvironment::WorkgroupSize() const {
-    return workgroup_size;
-}
-
-std::optional<Shader::ReplaceConstant> FileEnvironment::GetReplaceConstBuffer(u32 bank, u32 offset) {
-    auto const it = cbuf_replacements.find((u64(bank) << 32) | u64(offset));
-    return it != cbuf_replacements.end() ? std::optional{it->second} : std::nullopt;
-}
-
-int ShaderRecompilerImpl(int argc, char *argv[]) {
-    if (argc != 2) {
-        printf("usage: %s [input file] [-n]\n"
-            "RAW SPIRV CODE WILL BE SENT TO STDOUT!\n", argv[0]);
-        return EXIT_FAILURE;
-    }
-
-    size_t cfg_offset = 0;
-
-    Shader::ObjectPool<Shader::IR::Inst> inst_pool;
-    Shader::ObjectPool<Shader::IR::Block> block_pool;
-    Shader::ObjectPool<Shader::Maxwell::Flow::Block> cfg_blocks;
-    FileEnvironment env;
-
-    FILE *fp = fopen(argv[1], "rb");
-    if (fp != NULL) {
-        struct stat st;
-        fstat(fileno(fp), &st);
-        auto const words = (st.st_size / sizeof(u64));
-        env.code.resize(words + 1);
-        fread(env.code.data(), sizeof(u64), words, fp);
-        fclose(fp);
-    }
-
-    env.read_highest = env.read_lowest + env.code.size() * sizeof(u64);
-
-    Shader::Maxwell::Flow::CFG cfg(env, cfg_blocks, cfg_offset);
-
-    Shader::HostTranslateInfo host_info;
-    host_info.support_float64 = true;
-    host_info.support_float16 = true;
-    host_info.support_int64 = true;
-    host_info.needs_demote_reorder = true;
-    host_info.support_snorm_render_buffer = true;
-    host_info.support_viewport_index_layer = true;
-    host_info.support_geometry_shader_passthrough = true;
-    host_info.support_conditional_barrier = true;
-    host_info.min_ssbo_alignment = 0;
-    auto program = Shader::Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg, host_info);
-
-    // IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Block>& block_pool,
-    //                              Environment& env, Flow::CFG& cfg, const HostTranslateInfo& host_info)
-    // std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
-    //                            IR::Program& program, Bindings& bindings, bool optimize)
-    Shader::Profile profile{};
-    Shader::RuntimeInfo runtime_info;
-    auto const spirv_pgm = Shader::Backend::SPIRV::EmitSPIRV(profile, program, true);
-    fwrite(spirv_pgm.data(), sizeof(u64), spirv_pgm.size(), stdout);
-
-    return EXIT_SUCCESS;
-}

tools/maxwell-spirv/spirv_recompiler_impl.cpp

@@ -0,0 +1,67 @@
+// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
+// SPDX-License-Identifier: GPL-3.0-or-later
+
+#include <cstdlib>
+#include <sys/stat.h>
+#include "shader_recompiler/backend/spirv/emit_spirv.h"
+#include "shader_recompiler/environment.h"
+#include "shader_recompiler/frontend/maxwell/control_flow.h"
+#include "shader_recompiler/frontend/maxwell/translate_program.h"
+#include "shader_recompiler/host_translate_info.h"
+#include "shader_recompiler/object_pool.h"
+#include "shader_recompiler/profile.h"
+#include "shader_recompiler/runtime_info.h"
+
+#include "../maxwell-disas/file_environment.h"
+
+int SpirvShaderRecompilerImpl(int argc, char *argv[]) {
+    if (argc != 2) {
+        printf("usage: %s [input file] [-n]\n"
+            "RAW SPIRV CODE WILL BE SENT TO STDOUT!\n", argv[0]);
+        return EXIT_FAILURE;
+    }
+
+    size_t cfg_offset = 0;
+
+    Shader::ObjectPool<Shader::IR::Inst> inst_pool;
+    Shader::ObjectPool<Shader::IR::Block> block_pool;
+    Shader::ObjectPool<Shader::Maxwell::Flow::Block> cfg_blocks;
+    FileEnvironment env;
+
+    FILE *fp = fopen(argv[1], "rb");
+    if (fp != NULL) {
+        struct stat st;
+        fstat(fileno(fp), &st);
+        auto const words = (st.st_size / sizeof(u64));
+        env.code.resize(words + 1);
+        fread(env.code.data(), sizeof(u64), words, fp);
+        fclose(fp);
+    }
+
+    env.read_highest = env.read_lowest + env.code.size() * sizeof(u64);
+
+    Shader::Maxwell::Flow::CFG cfg(env, cfg_blocks, cfg_offset);
+
+    Shader::HostTranslateInfo host_info;
+    host_info.support_float64 = true;
+    host_info.support_float16 = true;
+    host_info.support_int64 = true;
+    host_info.needs_demote_reorder = true;
+    host_info.support_snorm_render_buffer = true;
+    host_info.support_viewport_index_layer = true;
+    host_info.support_geometry_shader_passthrough = true;
+    host_info.support_conditional_barrier = true;
+    host_info.min_ssbo_alignment = 0;
+    auto program = Shader::Maxwell::TranslateProgram(inst_pool, block_pool, env, cfg, host_info);
+
+    // IR::Program TranslateProgram(ObjectPool<IR::Inst>& inst_pool, ObjectPool<IR::Block>& block_pool,
+    //                              Environment& env, Flow::CFG& cfg, const HostTranslateInfo& host_info)
+    // std::vector<u32> EmitSPIRV(const Profile& profile, const RuntimeInfo& runtime_info,
+    //                            IR::Program& program, Bindings& bindings, bool optimize)
+    Shader::Profile profile{};
+    Shader::RuntimeInfo runtime_info;
+    auto const spirv_pgm = Shader::Backend::SPIRV::EmitSPIRV(profile, program, true);
+    fwrite(spirv_pgm.data(), sizeof(u64), spirv_pgm.size(), stdout);
+
+    return EXIT_SUCCESS;
+}