camerin
/
eden
mirror of https://git.eden-emu.dev/eden-emu/eden.git
1
0
Fork 0
Code Issues Packages Projects Releases Wiki Activity

[tools] add separate maxwell disassembler and spirv translator 

Signed-off-by: lizzie <lizzie@eden-emu.dev>
This commit is contained in:
lizzie 2026-02-02 06:26:32 +00:00
parent d0a054270e
commit b8805bc561
9 changed files with 1087 additions and 1 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

5
CMakeLists.txt
View File

@ -779,6 +779,11 @@ endif()
add_subdirectory(src)
if (TRUE OR ENABLE_TEGRA_TOOLS)
add_subdirectory(tools/maxwell-disas)
add_subdirectory(tools/maxwell-spirv)
endif()
# Set yuzu project or yuzu-cmd project as default StartUp Project in Visual Studio depending on whether QT is enabled or not
if(ENABLE_QT)
set_property(DIRECTORY ${CMAKE_CURRENT_SOURCE_DIR} PROPERTY VS_STARTUP_PROJECT yuzu)

7
tools/README.md
View File

@ -6,7 +6,12 @@ Tools for Eden and other subprojects.
- [CPMUtil Scripts](./cpm)
## Eden
## Binaries
- `maxwell-spirv`: Converts Maxwell shaders (dumped from `.ash` files) into SPIR-V code (emitted into STDOUT).
- `maxwell-disas`: Dumb maxwell dissasembler.
## Scripts
- `generate_converters.py`: Generates converters for given formats of textures (C++ helper).
- `svc_generator.py`: Generates the files `src/core/hle/kernel/svc.cpp` and `src/core/hle/kernel/svc.h` based off prototypes.

9
tools/maxwell-disas/CMakeLists.txt Normal file
View File

@ -0,0 +1,9 @@
# SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project
# SPDX-License-Identifier: GPL-3.0-or-later
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)
install(TARGETS maxwell-disas)
endif()
create_target_directory_groups(maxwell-disas)

638
tools/maxwell-disas/generated.cpp Normal file
View File

@ -0,0 +1,638 @@
#include <cstdint>
#include <string>
#include <string_view>
namespace Shader::Maxwell {
std::string DissasemblyFormat(uint64_t inst) {
std::string s{};
if(((inst>>48)&0xfff8ULL)==0x5c58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSWZADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c68ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c68ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RRO "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RRO "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5080ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MUFU "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c88ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCHK "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c88ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCHK "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c38ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c38ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c18ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISCADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c18ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISCADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c20ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c20ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFE "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFE "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5bf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFI "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4bf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFI "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x53f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFI "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c28ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c28ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5bf8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cf8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5be0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5be0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c30ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FLO "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c30ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FLO "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c08ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"POPC "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c08ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"POPC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cb8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ce0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ce0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ce0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ce0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5c98ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MOV "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4c98ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MOV "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ca0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SEL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ca0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SEL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xef10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHFL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xef10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHFL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xef10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHFL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xef10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHFL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5ce8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"P2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4ce8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"P2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5cf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"R2P "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4cf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"R2P "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5098ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CSET "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5098ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CSET "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x50a0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CSETP ";
if(((inst>>48)&0xfff8ULL)==0x50a0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CSETP ";
if(((inst>>48)&0xfff8ULL)==0x5088ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PSET "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5088ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PSET "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5090ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PSETP ";
if(((inst>>48)&0xfff8ULL)==0x5090ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PSETP ";
if(((inst>>48)&0xfff8ULL)==0xeea0ULL)return s+"STP ";
if(((inst>>48)&0xfff8ULL)==0xdf58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TMML "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TMML "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TMML "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TMML "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXQ "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xdf40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DEPBAR ";
if(((inst>>48)&0xfff8ULL)==0xf0f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DEPBAR ";
if(((inst>>48)&0xfff8ULL)==0xf0f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DEPBAR ";
if(((inst>>48)&0xfff8ULL)==0xefa0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AL2P "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefa0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AL2P "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xefd8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ALD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefd8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ALD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefd8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ALD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefd8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ALD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeff0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AST ";
if(((inst>>48)&0xfff8ULL)==0xeff0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AST ";
if(((inst>>48)&0xfff8ULL)==0xeff0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AST ";
if(((inst>>48)&0xfff8ULL)==0xeff0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"AST ";
if(((inst>>48)&0xfff8ULL)==0xfbe0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"OUT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xebe0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"OUT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xefe8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PIXLD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefe8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PIXLD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefe8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PIXLD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xef90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef90ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeed0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeed0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeec8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeec8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDL "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDL "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LDS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x5bd0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x4bd0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x5bd8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xeed8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xeed8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STG "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STL "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STL "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xef58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"STS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xebf8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RED ";
if(((inst>>48)&0xfff8ULL)==0xebf8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RED ";
if(((inst>>48)&0xfff8ULL)==0xef80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLL ";
if(((inst>>48)&0xfff8ULL)==0xef80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLL ";
if(((inst>>48)&0xfff8ULL)==0xef80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLL "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xef80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLL "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xebe8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLT ";
if(((inst>>48)&0xfff8ULL)==0xebf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLT ";
if(((inst>>48)&0xfff8ULL)==0xebf0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTLT ";
if(((inst>>48)&0xfff8ULL)==0xef98ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MEMBAR ";
if(((inst>>48)&0xfff8ULL)==0xeb10ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SULD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb18ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SULD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SULD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb08ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SULD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb30ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUST "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb38ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUST "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb20ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUST "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb28ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUST "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SURED "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb58ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SURED "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SURED "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xeb48ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SURED "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xea70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xea60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xea68ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xead0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xeac0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xeac8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0f8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SYNC ";
if(((inst>>48)&0xfff8ULL)==0x50b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"NOP ";
if(((inst>>48)&0xfff8ULL)==0x50b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"NOP ";
if(((inst>>48)&0xfff8ULL)==0xf0c8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"S2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+std::string(SpecialRegGetName((inst&0xfffffff)>>0x14));
if(((inst>>48)&0xfff8ULL)==0x50c8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CS2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+std::string(SpecialRegGetName((inst&0xfffffff)>>0x14));
if(((inst>>48)&0xfff8ULL)==0xf0b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"B2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"B2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"B2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0c0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"R2B "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x50d0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEPC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0xf0a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BAR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff8ULL)==0x50e0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VOTE ";
if(((inst>>48)&0xfff8ULL)==0x50d8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VOTE "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0xefd0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISBERD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff8ULL)==0x0ULL)return s+"HFMA2 ";
if(((inst>>48)&0xfff8ULL)==0x0ULL)return s+"HSET2 ";
if(((inst>>48)&0xfff8ULL)==0x0ULL)return s+"HSET2 ";
if(((inst>>48)&0xfff8ULL)==0x0ULL)return s+"HADD2 ";
if(((inst>>48)&0xfff8ULL)==0x0ULL)return s+"HMUL2 ";
if(((inst>>48)&0xfff8ULL)==0x5d20ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff8ULL)==0x5d20ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfef8ULL)==0x3858ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5ba0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4ba0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x53a0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3868ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3860ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5bb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5bb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4bb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4bb0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3890ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RRO "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x3888ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCHK "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5370ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3870ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3880ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3850ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3838ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMUL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3810ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5cc0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4cc0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3818ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISCADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3820ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFE "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x36f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BFI "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3828ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3848ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x36f8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x38f8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SHF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3840ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3830ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FLO "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b50ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5340ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5b40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4b40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5340ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x3808ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"POPC "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"F2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38b8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2F "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38e0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38e0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"I2I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x3898ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MOV "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0x100ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"MOV32I "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef8ULL)==0x38a0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SEL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x5bc0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PRMT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0x4bc0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PRMT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0x53c0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PRMT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x38e8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"P2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x38e8ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"P2R "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x38f0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"R2P "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0xf6e0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"OUT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef8ULL)==0x36d0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xeef0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xee60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xee70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0xee70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xee70ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfff0ULL)==0xe240ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BRA ";
if(((inst>>48)&0xfff0ULL)==0xe240ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BRA ";
if(((inst>>48)&0xfff0ULL)==0xe250ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BRX "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe250ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BRX "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe210ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"JMP ";
if(((inst>>48)&0xfff0ULL)==0xe210ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"JMP ";
if(((inst>>48)&0xfff0ULL)==0xe200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"JMX "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"JMX "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe260ULL)return s+"CAL ";
if(((inst>>48)&0xfff0ULL)==0xe260ULL)return s+"CAL ";
if(((inst>>48)&0xfff0ULL)==0xe270ULL)return s+"PRET ";
if(((inst>>48)&0xfff0ULL)==0xe270ULL)return s+"PRET ";
if(((inst>>48)&0xfff0ULL)==0xe220ULL)return s+"JCAL ";
if(((inst>>48)&0xfff0ULL)==0xe220ULL)return s+"JCAL ";
if(((inst>>48)&0xfff0ULL)==0xe290ULL)return s+"SSY ";
if(((inst>>48)&0xfff0ULL)==0xe290ULL)return s+"SSY ";
if(((inst>>48)&0xfff0ULL)==0xe280ULL)return s+"PLONGJMP ";
if(((inst>>48)&0xfff0ULL)==0xe280ULL)return s+"PLONGJMP ";
if(((inst>>48)&0xfff0ULL)==0xe2a0ULL)return s+"PBK ";
if(((inst>>48)&0xfff0ULL)==0xe2a0ULL)return s+"PBK ";
if(((inst>>48)&0xfff0ULL)==0xe2b0ULL)return s+"PCNT ";
if(((inst>>48)&0xfff0ULL)==0xe2b0ULL)return s+"PCNT ";
if(((inst>>48)&0xfff0ULL)==0xe320ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"RET ";
if(((inst>>48)&0xfff0ULL)==0xe310ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LONGJMP ";
if(((inst>>48)&0xfff0ULL)==0xe330ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"KIL ";
if(((inst>>48)&0xfff0ULL)==0xe340ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"BRK ";
if(((inst>>48)&0xfff0ULL)==0xe350ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CONT ";
if(((inst>>48)&0xfff0ULL)==0xe300ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"EXIT ";
if(((inst>>48)&0xfff0ULL)==0xe230ULL)return s+"PEXIT ";
if(((inst>>48)&0xfff0ULL)==0xe370ULL)return s+"SAM ";
if(((inst>>48)&0xfff0ULL)==0xe380ULL)return s+"RAM ";
if(((inst>>48)&0xfff0ULL)==0xe3a0ULL)return s+"BPT ";
if(((inst>>48)&0xfff0ULL)==0xe360ULL)return s+"RTT ";
if(((inst>>48)&0xfff0ULL)==0xe390ULL)return s+"IDE ";
if(((inst>>48)&0xfff0ULL)==0xe390ULL)return s+"IDE ";
if(((inst>>48)&0xfff0ULL)==0xe2e0ULL)return s+"SETCRSPTR "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe2c0ULL)return s+"GETCRSPTR "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfff0ULL)==0xe2f0ULL)return s+"SETLMEMBASE "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfff0ULL)==0xe2d0ULL)return s+"GETLMEMBASE "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xfef0ULL)==0x36a0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FCMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x36b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x36b0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3670ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3680ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3680ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x38c0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3650ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3650ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3650ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3650ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3660ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3660ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3660ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3660ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISETP "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3640ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x3640ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ICMP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfef0ULL)==0x36c0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"PRMT "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffe0ULL)==0xee40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL ";
if(((inst>>48)&0xffe0ULL)==0xef60ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"CCTL ";
if(((inst>>48)&0xffc0ULL)==0x5b00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"XMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xdf00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4S "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xdf80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4S "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xdec0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xdec0ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xde40ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TXD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xee00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xeb80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xea00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xffc0ULL)==0xea80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"SUATOM "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x5980ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x4980ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5180ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5900ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x5900ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x4900ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x4900ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x4a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5a80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMADSP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x4a80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMADSP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5280ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMADSP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfec0ULL)==0x3600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"XMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfec0ULL)==0x3600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"XMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5100ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"XMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff80ULL)==0x5000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x5000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff80ULL)==0x5000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe80ULL)==0x3280ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x1e00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FMUL32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x4800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x4800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0xe000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IPA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xe000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IPA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xe000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IPA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xe000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IPA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe80ULL)==0x3200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x3200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"DSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x3400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x3480ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMADSP "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x1f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMUL32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5f00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VABSDIFF "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5700ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHL "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0x5600ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSHR "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xff00ULL)==0xdc00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xdd00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xed00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xed00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOM "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xec00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xff00ULL)==0xec00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ATOMS "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HSET2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HSET2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HSET2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HSET2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HADD2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HADD2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HMUL2 ";
if(((inst>>48)&0xfe80ULL)==0x80ULL)return s+"HMUL2 ";
if(((inst>>48)&0xfe80ULL)==0x7e00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x7e00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x7e80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe80ULL)==0x7e80ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HSETP2 "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x1c00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IADD32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4e00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"XMAD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x3a00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VMNMX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x4000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VSET "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0xd800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEXS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xd800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEXS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xd000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEXS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xd000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEXS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xda00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLDS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xda00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLDS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xd200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLDS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0xd200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLDS "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfe00ULL)==0x0ULL)return s+"HMUL2_32I ";
if(((inst>>48)&0xfe00ULL)==0x0ULL)return s+"HADD2_32I ";
if(((inst>>48)&0xfe00ULL)==0x2800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HFMA2_32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x2800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HFMA2_32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0xc00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0xc00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FFMA32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"FADD32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x1000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"IMAD32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x1400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ISCADD32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP32I "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x400ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfe00ULL)==0x200ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x3c00ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LOP3 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x2000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"VADD "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xfc00ULL)==0x1800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LEA "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xf880ULL)==0x0ULL)return s+"HFMA2 ";
if(((inst>>48)&0xf880ULL)==0x0ULL)return s+"HFMA2 ";
if(((inst>>48)&0xf880ULL)==0x6080ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"HFMA2 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" ";
if(((inst>>48)&0xf800ULL)==0xc000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xf800ULL)==0xc000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TEX "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xf800ULL)==0xc800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xf800ULL)==0xc800ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"TLD4 "+"R"+std::to_string((inst&0xff)>>0x0)+" "+"R"+std::to_string((inst&0xffff)>>0x8)+" "+"R"+std::to_string((inst&0xfffffff)>>0x14)+" ";
if(((inst>>48)&0xe000ULL)==0x8000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xe000ULL)==0x8000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"LD "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xe000ULL)==0xa000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ST "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
if(((inst>>48)&0xe000ULL)==0xa000ULL)return s+(((inst&0x7ffff)>>0x10)?"":"!")+"ST "+"R"+std::to_string((inst&0xff)>>0x0)+" ";
return "?";}
}

219
tools/maxwell-disas/main.cpp Normal file
View File

@ -0,0 +1,219 @@
// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
// SPDX-License-Identifier: GPL-3.0-or-later
#include "common/assert.h"
#include <cstdlib>
#include <sys/stat.h>
#include <vector>
enum class Opcode {
#define INST(name, cute, encode) name,
#include "shader_recompiler/frontend/maxwell/maxwell.inc"
#undef INST
};
consteval std::pair<u64, u64> MaskValueFromEncoding(const char data[20]) noexcept {
u64 mask = 0, value = 0, bit = u64(1) << 63;
for (int i = 0; i < 20; ++i)
switch (data[i]) {
case '0':
mask |= bit;
bit >>= 1;
break;
case '1':
mask |= bit;
value |= bit;
bit >>= 1;
break;
case '-':
bit >>= 1;
break;
default:
break;
}
return { mask, value };
}
Opcode Decode(u64 insn) {
#define INST(name, cute, encode) \
if (auto const p = MaskValueFromEncoding(encode); (insn & p.first) == p.second) \
return Opcode::name;
#include "shader_recompiler/frontend/maxwell/maxwell.inc"
#undef INST
ASSERT_MSG(false, "Invalid insn 0x{:016x}", insn);
return Opcode::NOP;
}
const char* NameOf(Opcode opcode) {
constexpr const char* NAME_TABLE[] = {
#define INST(name, cute, encode) cute,
#include "shader_recompiler/frontend/maxwell/maxwell.inc"
#undef INST
};
ASSERT_MSG(size_t(opcode) < sizeof(NAME_TABLE) / sizeof(NAME_TABLE[0]), "Invalid opcode with raw value {}", int(opcode));
return NAME_TABLE[size_t(opcode)];
}
namespace Shader::Maxwell {
std::string_view SpecialRegGetName(size_t i) {
switch (i) {
case 0: return "SR_LANEID";
case 1: return "SR_CLOCK";
case 2: return "SR_VIRTCFG";
case 3: return "SR_VIRTID";
case 4: return "SR_PM0";
case 5: return "SR_PM1";
case 6: return "SR_PM2";
case 7: return "SR_PM3";
case 8: return "SR_PM4";
case 9: return "SR_PM5";
case 10: return "SR_PM6";
case 11: return "SR_PM7";
case 12: return "SR_?";
case 13: return "SR_?";
case 14: return "SR_?";
case 15: return "SR_ORDERING_TICKET";
case 16: return "SR_PRIM_TYPE";
case 17: return "SR_INVOCATION_ID";
case 18: return "SR_Y_DIRECTION";
case 19: return "SR_THREAD_KILL";
case 20: return "SM_SHADER_TYPE";
case 21: return "SR_DIRECTCBEWRITEADDRESSLOW";
case 22: return "SR_DIRECTCBEWRITEADDRESSHIGH";
case 23: return "SR_DIRECTCBEWRITEENABLED";
case 24: return "SR_MACHINE_ID_0";
case 25: return "SR_MACHINE_ID_1";
case 26: return "SR_MACHINE_ID_2";
case 27: return "SR_MACHINE_ID_3";
case 28: return "SR_AFFINITY";
case 29: return "SR_INVOCATION_INFO";
case 30: return "SR_WSCALEFACTOR_XY";
case 31: return "SR_WSCALEFACTOR_Z";
case 32: return "SR_TID";
case 33: return "SR_TID_X";
case 34: return "SR_TID_Y";
case 35: return "SR_TID_Z";
case 36: return "SR_CTA_PARAM";
case 37: return "SR_CTAID_X";
case 38: return "SR_CTAID_Y";
case 39: return "SR_CTAID_Z";
case 40: return "SR_NTID";
case 41: return "SR_CirQueueIncrMinusOne";
case 42: return "SR_NLATC";
case 43: return "SR_?";
case 44: return "SR_SM_SPA_VERSION";
case 45: return "SR_MULTIPASSSHADERINFO";
case 46: return "SR_LWINHI";
case 47: return "SR_SWINHI";
case 48: return "SR_SWINLO";
case 49: return "SR_SWINSZ";
case 50: return "SR_SMEMSZ";
case 51: return "SR_SMEMBANKS";
case 52: return "SR_LWINLO";
case 53: return "SR_LWINSZ";
case 54: return "SR_LMEMLOSZ";
case 55: return "SR_LMEMHIOFF";
case 56: return "SR_EQMASK";
case 57: return "SR_LTMASK";
case 58: return "SR_LEMASK";
case 59: return "SR_GTMASK";
case 60: return "SR_GEMASK";
case 61: return "SR_REGALLOC";
case 62: return "SR_BARRIERALLOC";
case 63: return "SR_?";
case 64: return "SR_GLOBALERRORSTATUS";
case 65: return "SR_?";
case 66: return "SR_WARPERRORSTATUS";
case 67: return "SR_WARPERRORSTATUSCLEAR";
case 68: return "SR_?";
case 69: return "SR_?";
case 70: return "SR_?";
case 71: return "SR_?";
case 72: return "SR_PM_HI0";
case 73: return "SR_PM_HI1";
case 74: return "SR_PM_HI2";
case 75: return "SR_PM_HI3";
case 76: return "SR_PM_HI4";
case 77: return "SR_PM_HI5";
case 78: return "SR_PM_HI6";
case 79: return "SR_PM_HI7";
case 80: return "SR_CLOCKLO";
case 81: return "SR_CLOCKHI";
case 82: return "SR_GLOBALTIMERLO";
case 83: return "SR_GLOBALTIMERHI";
case 84: return "SR_?";
case 85: return "SR_?";
case 86: return "SR_?";
case 87: return "SR_?";
case 88: return "SR_?";
case 89: return "SR_?";
case 90: return "SR_?";
case 91: return "SR_?";
case 92: return "SR_?";
case 93: return "SR_?";
case 94: return "SR_?";
case 95: return "SR_?";
case 96: return "SR_HWTASKID";
case 97: return "SR_CIRCULARQUEUEENTRYINDEX";
case 98: return "SR_CIRCULARQUEUEENTRYADDRESSLOW";
case 99: return "SR_CIRCULARQUEUEENTRYADDRESSHIGH";
default: return "SR_??"; }
}
}
#include "generated.cpp"
int ReferenceImpl(int argc, char *argv[]) {
std::vector<uint64_t> code;
FILE *fp = fopen(argv[1], "rb");
if (fp != NULL) {
struct stat st;
fstat(fileno(fp), &st);
auto const words = (size_t(st.st_size) / sizeof(uint64_t));
code.resize(words + 1);
fread(code.data(), sizeof(uint64_t), words, fp);
fclose(fp);
}
for (size_t i = 0; i < code.size(); ++i) {
printf("%016lx\t%-40s\n", code[i]
, Shader::Maxwell::DissasemblyFormat(code[i]).data()
);
}
return EXIT_SUCCESS;
}
int ShaderRecompilerDisas(int argc, char *argv[]) {
std::vector<u64> code;
FILE *fp = fopen(argv[1], "rb");
if (fp != NULL) {
struct stat st;
fstat(fileno(fp), &st);
auto const words = (size_t(st.st_size) / sizeof(u64));
code.resize(words + 1);
fread(code.data(), sizeof(u64), words, fp);
fclose(fp);
}
for (size_t i = 0; i < code.size(); ++i) {
auto const opcode = Decode(code[i]);
printf("%016lx\t%s\n", code[i], NameOf(opcode));
}
return EXIT_SUCCESS;
}
int main(int argc, char *argv[]) {
if (argc < 2) {
printf(
"usage: %s [input file] [-n]\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;
}
if (argc >= 3) {
if (::strcmp(argv[2], "-n") == 0
|| ::strcmp(argv[2], "--new") == 0) {
return ReferenceImpl(argc, argv);
}
}
return ShaderRecompilerDisas(argc, argv);
}

13
tools/maxwell-spirv/CMakeLists.txt Normal file
View File

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

26
tools/maxwell-spirv/main.cpp Normal file
View File

@ -0,0 +1,26 @@
// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
// SPDX-License-Identifier: GPL-3.0-or-later
#include <cstdio>
#include <cstdlib>
#include <cstring>
int ReferenceImpl(int argc, char *argv[]);
int ShaderRecompilerImpl(int argc, char *argv[]);
int main(int argc, char *argv[]) {
if (argc < 2) {
printf("usage: %s [input file]\n"
"Specify -n to use a recompiler that is NOT tied to the shader recompiler\n"
"aka. a reference recompiler\n"
"RAW SPIRV CODE WILL BE SENT TO STDOUT!\n", argv[0]);
return EXIT_FAILURE;
}
if (argc >= 3) {
if (::strcmp(argv[2], "-n") == 0
|| ::strcmp(argv[2], "--new") == 0) {
return ReferenceImpl(argc, argv);
}
}
return ShaderRecompilerImpl(argc, argv);
}

162
tools/maxwell-spirv/recompiler_impl.cpp Normal file
View File

@ -0,0 +1,162 @@
// 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;
}

9
tools/maxwell-spirv/reference_impl.cpp Normal file
View File

@ -0,0 +1,9 @@
#include <cassert>
#include <cstdlib>
#include <sirit/sirit.h>
#include <sys/stat.h>
int ReferenceImpl(int argc, char *argv[]) {
//todo
return EXIT_SUCCESS;
}