diff --git a/CMakeLists.txt b/CMakeLists.txt index de57f10e70..3006772c6e 100644 --- a/CMakeLists.txt +++ b/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 diff --git a/tools/README.md b/tools/README.md index 93c74de1b7..4a633a4d3e 100644 --- a/tools/README.md +++ b/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 diff --git a/tools/maxwell-disas/CMakeLists.txt b/tools/maxwell-disas/CMakeLists.txt index 194cbe6a39..2595d45a35 100644 --- a/tools/maxwell-disas/CMakeLists.txt +++ b/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) diff --git a/tools/maxwell-disas/file_environment.h b/tools/maxwell-disas/file_environment.h new file mode 100644 index 0000000000..cf190806e4 --- /dev/null +++ b/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 +#include +#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 WorkgroupSize() const override { return workgroup_size; } + [[nodiscard]] std::optional 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 code; + std::unordered_map texture_types; + std::unordered_map texture_pixel_formats; + std::unordered_map cbuf_values; + std::unordered_map cbuf_replacements; + std::array 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; +}; diff --git a/tools/maxwell-disas/main.cpp b/tools/maxwell-disas/main.cpp index 6e20855b42..0feee75824 100644 --- a/tools/maxwell-disas/main.cpp +++ b/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 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 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 - || ::strcmp(argv[2], "--new") == 0) { - return ReferenceImpl(argc, argv); + if (::strcmp(argv[2], "-n") == 0 || ::strcmp(argv[2], "--new") == 0) { + return DisasReferenceImpl(argc, argv); } } - return ShaderRecompilerDisas(argc, argv); + return DisasShaderRecompilerImpl(argc, argv); } diff --git a/tools/maxwell-ir/CMakeLists.txt b/tools/maxwell-ir/CMakeLists.txt new file mode 100644 index 0000000000..45441dcae1 --- /dev/null +++ b/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) diff --git a/tools/maxwell-ir/main.cpp b/tools/maxwell-ir/main.cpp new file mode 100644 index 0000000000..66aeaeac40 --- /dev/null +++ b/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 +#include +#include + +#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 inst_pool; + Shader::ObjectPool block_pool; + Shader::ObjectPool 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); +} diff --git a/tools/maxwell-spirv/CMakeLists.txt b/tools/maxwell-spirv/CMakeLists.txt index 323d5046f8..819373712e 100644 --- a/tools/maxwell-spirv/CMakeLists.txt +++ b/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 - reference_impl.cpp + spirv_recompiler_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) diff --git a/tools/maxwell-spirv/main.cpp b/tools/maxwell-spirv/main.cpp index 46d7d62202..f9245d436e 100644 --- a/tools/maxwell-spirv/main.cpp +++ b/tools/maxwell-spirv/main.cpp @@ -5,8 +5,8 @@ #include #include -int ReferenceImpl(int argc, char *argv[]); -int ShaderRecompilerImpl(int argc, char *argv[]); +int SpirvReferenceImpl(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); } diff --git a/tools/maxwell-spirv/recompiler_impl.cpp b/tools/maxwell-spirv/recompiler_impl.cpp deleted file mode 100644 index 4666d5b9fe..0000000000 --- a/tools/maxwell-spirv/recompiler_impl.cpp +++ /dev/null @@ -1,162 +0,0 @@ -// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project -// SPDX-License-Identifier: GPL-3.0-or-later - -#include -#include -#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 WorkgroupSize() const override; - [[nodiscard]] std::optional 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 code; - std::unordered_map texture_types; - std::unordered_map texture_pixel_formats; - std::unordered_map cbuf_values; - std::unordered_map cbuf_replacements; - std::array 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 FileEnvironment::WorkgroupSize() const { - return workgroup_size; -} - -std::optional 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 inst_pool; - Shader::ObjectPool block_pool; - Shader::ObjectPool 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& inst_pool, ObjectPool& block_pool, - // Environment& env, Flow::CFG& cfg, const HostTranslateInfo& host_info) - // std::vector 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; -} diff --git a/tools/maxwell-spirv/spirv_recompiler_impl.cpp b/tools/maxwell-spirv/spirv_recompiler_impl.cpp new file mode 100644 index 0000000000..55830abe39 --- /dev/null +++ b/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 +#include +#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 inst_pool; + Shader::ObjectPool block_pool; + Shader::ObjectPool 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& inst_pool, ObjectPool& block_pool, + // Environment& env, Flow::CFG& cfg, const HostTranslateInfo& host_info) + // std::vector 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; +} diff --git a/tools/maxwell-spirv/reference_impl.cpp b/tools/maxwell-spirv/spirv_reference_impl.cpp similarity index 100% rename from tools/maxwell-spirv/reference_impl.cpp rename to tools/maxwell-spirv/spirv_reference_impl.cpp