diff --git a/CMakeLists.txt b/CMakeLists.txt index 22a811d30..95766bc67 100755 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -766,11 +766,11 @@ set(SHADER_RECOMPILER src/shader_recompiler/exception.h src/shader_recompiler/ir/passes/identity_removal_pass.cpp src/shader_recompiler/ir/passes/ir_passes.h src/shader_recompiler/ir/passes/lower_buffer_format_to_raw.cpp - src/shader_recompiler/ir/passes/lower_shared_mem_to_registers.cpp src/shader_recompiler/ir/passes/resource_tracking_pass.cpp src/shader_recompiler/ir/passes/ring_access_elimination.cpp src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp + src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp src/shader_recompiler/ir/passes/ssa_rewrite_pass.cpp src/shader_recompiler/ir/abstract_syntax_list.h src/shader_recompiler/ir/attribute.cpp diff --git a/src/common/config.cpp b/src/common/config.cpp index aae903da6..32c5e670b 100644 --- a/src/common/config.cpp +++ b/src/common/config.cpp @@ -68,6 +68,7 @@ static bool vkCrashDiagnostic = false; static bool vkHostMarkers = false; static bool vkGuestMarkers = false; static bool rdocEnable = false; +static bool isFpsColor = true; static s16 cursorState = HideCursorState::Idle; static int cursorHideTimeout = 5; // 5 seconds (default) static bool useUnifiedInputConfig = true; @@ -282,6 +283,10 @@ bool isRdocEnabled() { return rdocEnable; } +bool fpsColor() { + return isFpsColor; +} + u32 vblankDiv() { return vblankDivider; } @@ -757,6 +762,7 @@ void load(const std::filesystem::path& path) { isDebugDump = toml::find_or(debug, "DebugDump", false); isShaderDebug = toml::find_or(debug, "CollectShader", false); + isFpsColor = toml::find_or(debug, "FPSColor", true); } if (data.contains("GUI")) { @@ -807,8 +813,8 @@ void load(const std::filesystem::path& path) { // Check if the loaded language is in the allowed list const std::vector allowed_languages = { "ar_SA", "da_DK", "de_DE", "el_GR", "en_US", "es_ES", "fa_IR", "fi_FI", "fr_FR", "hu_HU", - "id_ID", "it_IT", "ja_JP", "ko_KR", "lt_LT", "nl_NL", "no_NO", "pl_PL", "pt_BR", "ro_RO", - "ru_RU", "sq_AL", "sv_SE", "tr_TR", "uk_UA", "vi_VN", "zh_CN", "zh_TW"}; + "id_ID", "it_IT", "ja_JP", "ko_KR", "lt_LT", "nb_NO", "nl_NL", "pl_PL", "pt_BR", "pt_PT", + "ro_RO", "ru_RU", "sq_AL", "sv_SE", "tr_TR", "uk_UA", "vi_VN", "zh_CN", "zh_TW"}; if (std::find(allowed_languages.begin(), allowed_languages.end(), emulator_language) == allowed_languages.end()) { @@ -881,6 +887,7 @@ void save(const std::filesystem::path& path) { data["Vulkan"]["rdocEnable"] = rdocEnable; data["Debug"]["DebugDump"] = isDebugDump; data["Debug"]["CollectShader"] = isShaderDebug; + data["Debug"]["FPSColor"] = isFpsColor; data["Keys"]["TrophyKey"] = trophyKey; diff --git a/src/common/config.h b/src/common/config.h index dfb1d9fad..7b9bc789b 100644 --- a/src/common/config.h +++ b/src/common/config.h @@ -67,6 +67,7 @@ bool copyGPUCmdBuffers(); bool dumpShaders(); bool patchShaders(); bool isRdocEnabled(); +bool fpsColor(); u32 vblankDiv(); void setDebugDump(bool enable); diff --git a/src/core/devtools/layer.cpp b/src/core/devtools/layer.cpp index a6d99b49b..603d76df5 100644 --- a/src/core/devtools/layer.cpp +++ b/src/core/devtools/layer.cpp @@ -259,7 +259,19 @@ void L::DrawAdvanced() { void L::DrawSimple() { const float frameRate = DebugState.Framerate; + if (Config::fpsColor()) { + if (frameRate < 10) { + PushStyleColor(ImGuiCol_Text, ImVec4(1.0f, 0.0f, 0.0f, 1.0f)); // Red + } else if (frameRate >= 10 && frameRate < 20) { + PushStyleColor(ImGuiCol_Text, ImVec4(1.0f, 0.5f, 0.0f, 1.0f)); // Orange + } else { + PushStyleColor(ImGuiCol_Text, ImVec4(1.0f, 1.0f, 1.0f, 1.0f)); // White + } + } else { + PushStyleColor(ImGuiCol_Text, ImVec4(1.0f, 1.0f, 1.0f, 1.0f)); // White + } Text("%d FPS (%.1f ms)", static_cast(std::round(frameRate)), 1000.0f / frameRate); + PopStyleColor(); } static void LoadSettings(const char* line) { diff --git a/src/qt_gui/cheats_patches.cpp b/src/qt_gui/cheats_patches.cpp index 866ab3ca0..e9db88381 100644 --- a/src/qt_gui/cheats_patches.cpp +++ b/src/qt_gui/cheats_patches.cpp @@ -91,9 +91,11 @@ void CheatsPatches::setupUI() { gameVersionLabel->setAlignment(Qt::AlignLeft); gameInfoLayout->addWidget(gameVersionLabel); - QLabel* gameSizeLabel = new QLabel(tr("Size: ") + m_gameSize); - gameSizeLabel->setAlignment(Qt::AlignLeft); - gameInfoLayout->addWidget(gameSizeLabel); + if (m_gameSize.left(4) != "0.00") { + QLabel* gameSizeLabel = new QLabel(tr("Size: ") + m_gameSize); + gameSizeLabel->setAlignment(Qt::AlignLeft); + gameInfoLayout->addWidget(gameSizeLabel); + } // Add a text area for instructions and 'Patch' descriptions instructionsTextEdit = new QTextEdit(); diff --git a/src/qt_gui/control_settings.cpp b/src/qt_gui/control_settings.cpp index 644576feb..73622e6b0 100644 --- a/src/qt_gui/control_settings.cpp +++ b/src/qt_gui/control_settings.cpp @@ -3,9 +3,9 @@ #include #include +#include #include "common/path_util.h" #include "control_settings.h" -#include "kbm_config_dialog.h" #include "ui_control_settings.h" ControlSettings::ControlSettings(std::shared_ptr game_info_get, QWidget* parent) @@ -16,7 +16,7 @@ ControlSettings::ControlSettings(std::shared_ptr game_info_get, Q AddBoxItems(); SetUIValuestoMappings(); - ui->KBMButton->setFocus(); + UpdateLightbarColor(); connect(ui->buttonBox, &QDialogButtonBox::clicked, this, [this](QAbstractButton* button) { if (button == ui->buttonBox->button(QDialogButtonBox::Save)) { @@ -29,11 +29,7 @@ ControlSettings::ControlSettings(std::shared_ptr game_info_get, Q }); connect(ui->buttonBox, &QDialogButtonBox::rejected, this, &QWidget::close); - connect(ui->KBMButton, &QPushButton::clicked, this, [this] { - auto KBMWindow = new EditorDialog(this); - KBMWindow->exec(); - SetUIValuestoMappings(); - }); + connect(ui->ProfileComboBox, &QComboBox::currentTextChanged, this, [this] { GetGameTitle(); SetUIValuestoMappings(); @@ -61,6 +57,27 @@ ControlSettings::ControlSettings(std::shared_ptr game_info_get, Q [this](int value) { ui->RStickLeftBox->setCurrentIndex(value); }); connect(ui->RStickLeftBox, &QComboBox::currentIndexChanged, this, [this](int value) { ui->RStickRightBox->setCurrentIndex(value); }); + + connect(ui->RSlider, &QSlider::valueChanged, this, [this](int value) { + QString RedValue = QString("%1").arg(value, 3, 10, QChar('0')); + QString RValue = "R: " + RedValue; + ui->RLabel->setText(RValue); + UpdateLightbarColor(); + }); + + connect(ui->GSlider, &QSlider::valueChanged, this, [this](int value) { + QString GreenValue = QString("%1").arg(value, 3, 10, QChar('0')); + QString GValue = "G: " + GreenValue; + ui->GLabel->setText(GValue); + UpdateLightbarColor(); + }); + + connect(ui->BSlider, &QSlider::valueChanged, this, [this](int value) { + QString BlueValue = QString("%1").arg(value, 3, 10, QChar('0')); + QString BValue = "B: " + BlueValue; + ui->BLabel->setText(BValue); + UpdateLightbarColor(); + }); } void ControlSettings::SaveControllerConfig(bool CloseOnSave) { @@ -121,7 +138,7 @@ void ControlSettings::SaveControllerConfig(bool CloseOnSave) { if (std::find(ControllerInputs.begin(), ControllerInputs.end(), input_string) != ControllerInputs.end() || - output_string == "analog_deadzone") { + output_string == "analog_deadzone" || output_string == "override_controller_color") { line.erase(); continue; } @@ -227,6 +244,14 @@ void ControlSettings::SaveControllerConfig(bool CloseOnSave) { deadzonevalue = std::to_string(ui->RightDeadzoneSlider->value()); lines.push_back("analog_deadzone = rightjoystick, " + deadzonevalue + ", 127"); + lines.push_back(""); + std::string OverrideLB = ui->LightbarCheckBox->isChecked() ? "true" : "false"; + std::string LightBarR = std::to_string(ui->RSlider->value()); + std::string LightBarG = std::to_string(ui->GSlider->value()); + std::string LightBarB = std::to_string(ui->BSlider->value()); + lines.push_back("override_controller_color = " + OverrideLB + ", " + LightBarR + ", " + + LightBarG + ", " + LightBarB); + std::vector save; bool CurrentLineEmpty = false, LastLineEmpty = false; for (auto const& line : lines) { @@ -243,6 +268,9 @@ void ControlSettings::SaveControllerConfig(bool CloseOnSave) { output_file.close(); Config::SetUseUnifiedInputConfig(!ui->PerGameCheckBox->isChecked()); + Config::SetOverrideControllerColor(ui->LightbarCheckBox->isChecked()); + Config::SetControllerCustomColor(ui->RSlider->value(), ui->GSlider->value(), + ui->BSlider->value()); Config::save(Common::FS::GetUserPath(Common::FS::PathType::UserDir) / "config.toml"); if (CloseOnSave) @@ -351,7 +379,7 @@ void ControlSettings::SetUIValuestoMappings() { if (std::find(ControllerInputs.begin(), ControllerInputs.end(), input_string) != ControllerInputs.end() || - output_string == "analog_deadzone") { + output_string == "analog_deadzone" || output_string == "override_controller_color") { if (input_string == "cross") { ui->ABox->setCurrentText(QString::fromStdString(output_string)); CrossExists = true; @@ -436,9 +464,45 @@ void ControlSettings::SetUIValuestoMappings() { ui->RightDeadzoneSlider->setValue(2); ui->RightDeadzoneValue->setText("2"); } + } else if (output_string == "override_controller_color") { + std::size_t comma_pos = line.find(','); + if (comma_pos != std::string::npos) { + std::string overridestring = line.substr(equal_pos + 1, comma_pos); + bool override = overridestring.contains("true") ? true : false; + ui->LightbarCheckBox->setChecked(override); + + std::string lightbarstring = line.substr(comma_pos + 1); + std::size_t comma_pos2 = lightbarstring.find(','); + if (comma_pos2 != std::string::npos) { + std::string Rstring = lightbarstring.substr(0, comma_pos2); + ui->RSlider->setValue(std::stoi(Rstring)); + QString RedValue = QString("%1").arg(std::stoi(Rstring), 3, 10, QChar('0')); + QString RValue = "R: " + RedValue; + ui->RLabel->setText(RValue); + } + + std::string GBstring = lightbarstring.substr(comma_pos2 + 1); + std::size_t comma_pos3 = GBstring.find(','); + if (comma_pos3 != std::string::npos) { + std::string Gstring = GBstring.substr(0, comma_pos3); + ui->GSlider->setValue(std::stoi(Gstring)); + QString GreenValue = + QString("%1").arg(std::stoi(Gstring), 3, 10, QChar('0')); + QString GValue = "G: " + GreenValue; + ui->GLabel->setText(GValue); + + std::string Bstring = GBstring.substr(comma_pos3 + 1); + ui->BSlider->setValue(std::stoi(Bstring)); + QString BlueValue = + QString("%1").arg(std::stoi(Bstring), 3, 10, QChar('0')); + QString BValue = "B: " + BlueValue; + ui->BLabel->setText(BValue); + } + } } } } + file.close(); // If an entry does not exist in the config file, we assume the user wants it unmapped if (!CrossExists) @@ -490,8 +554,6 @@ void ControlSettings::SetUIValuestoMappings() { ui->RStickUpBox->setCurrentText("unmapped"); ui->RStickDownBox->setCurrentText("unmapped"); } - - file.close(); } void ControlSettings::GetGameTitle() { @@ -507,4 +569,13 @@ void ControlSettings::GetGameTitle() { } } +void ControlSettings::UpdateLightbarColor() { + ui->LightbarColorFrame->setStyleSheet(""); + QString RValue = QString::number(ui->RSlider->value()); + QString GValue = QString::number(ui->GSlider->value()); + QString BValue = QString::number(ui->BSlider->value()); + QString colorstring = "background-color: rgb(" + RValue + "," + GValue + "," + BValue + ")"; + ui->LightbarColorFrame->setStyleSheet(colorstring); +} + ControlSettings::~ControlSettings() {} diff --git a/src/qt_gui/control_settings.h b/src/qt_gui/control_settings.h index 04227f3a8..e686f044d 100644 --- a/src/qt_gui/control_settings.h +++ b/src/qt_gui/control_settings.h @@ -18,6 +18,7 @@ public: private Q_SLOTS: void SaveControllerConfig(bool CloseOnSave); void SetDefault(); + void UpdateLightbarColor(); private: std::unique_ptr ui; diff --git a/src/qt_gui/control_settings.ui b/src/qt_gui/control_settings.ui index b6acb5ca9..e88e239e9 100644 --- a/src/qt_gui/control_settings.ui +++ b/src/qt_gui/control_settings.ui @@ -11,8 +11,8 @@ 0 0 - 1012 - 721 + 1043 + 792 @@ -25,760 +25,681 @@ - - QFrame::Shape::NoFrame - - - 0 - true - + 0 0 - 994 - 673 + 1019 + 732 - - - Control Settings - - - - 5 - - - 5 - - - 5 - - - 5 - + + + + 0 + 0 + 1021 + 731 + + + - + + + 5 + - - - 5 + + + true - - - - true - - - - 0 - 0 - - - - - 16777215 - 16777215 - - - - D-Pad - - - - 6 - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - 0 + + + 0 + 0 + + + + + 16777215 + 16777215 + + + + D-Pad + + + + 6 + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 124 + 0 + + + + 0 + 16777215 + + + + Up + + + + + + false + + + QComboBox::SizeAdjustPolicy::AdjustToContents + + + + + + + + + + + + + + + Left + + - 0 + 5 - 0 + 5 - 0 + 5 - 0 + 5 - - - - 124 - 0 - - - - - 0 - 16777215 - - - - Up - - - - - - false - - - QComboBox::SizeAdjustPolicy::AdjustToContents - - - - - + - - - - - Left - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - - - - Right - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - false - - - - - - - - - - - + + + Right + + - 0 + 5 - 0 + 5 - 0 + 5 - 0 + 5 - - - - 124 - 16777215 - + + + false - - Down - - - - - - true - - - - 0 - 0 - - - - - 0 - 0 - - - - - - - - - - - Qt::Orientation::Vertical - - - QSizePolicy::Policy::Maximum - - - - 20 - 40 - - - - - - - - - 0 - 0 - - - - Left Stick Deadzone (def:2 max:127) - - - - - - - 0 - 0 - - - - Left Deadzone - - - Qt::AlignmentFlag::AlignRight|Qt::AlignmentFlag::AlignTrailing|Qt::AlignmentFlag::AlignVCenter - - - - - - - - 0 - 0 - - - - 1 - - - 127 - - - Qt::Orientation::Horizontal - - - - - - - - - - - 0 - 0 - - - - - 0 - 0 - - - - Left Stick - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - 16777215 - 2121 - - - - - 0 + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 124 + 16777215 + - - 0 + + Down - - 0 - - - 0 - - - - - - 0 - 0 - - - - - 124 - 16777215 - - - - Up - - - - - - true - - - - - - - - - - - - - - - Left - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - true - - - - - - - - - - - 179 - 16777215 - - - - Right - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - true - - - - - - - - - - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 124 - 0 - - - - - 124 - 21212 - - - - Down - - - - - - true - - - false - - - false - - - - - - - - - - - - - + + + + + true + + + + 0 + 0 + + + + + 0 + 0 + + + + + + + + + + + + - - - 0 + + + Qt::Orientation::Vertical - - - - - 0 - 0 - - - - - 12 - true - - - - Config Selection - - - Qt::AlignmentFlag::AlignCenter - - + + QSizePolicy::Policy::Maximum + + + + 20 + 40 + + + + + + + + + 0 + 0 + + + + Left Stick Deadzone (def:2 max:127) + + + + + + + 0 + 0 + + + + Left Deadzone + + + Qt::AlignmentFlag::AlignRight|Qt::AlignmentFlag::AlignTrailing|Qt::AlignmentFlag::AlignVCenter + + + + + + + + 0 + 0 + + + + 1 + + + 127 + + + Qt::Orientation::Horizontal + + + + + + + + + + + 0 + 0 + + + + + 0 + 0 + + + + Left Stick + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + 16777215 + 2121 + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 0 + 0 + + + + + 124 + 16777215 + + + + Up + + + + + + true + + + + + + + + + + + - - - - - - 9 - false - - - - - - - -1 - - - Common Config - - - - - - - - 10 - true - - - - Common Config - - - Qt::AlignmentFlag::AlignCenter - - - true - - - - + + + Left + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + true + + + + + - - - - 0 - 0 - + + + + 179 + 16777215 + + + Right + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + true + + + + + + + + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 124 + 0 + + + + + 124 + 21212 + + + + Down + + + + + + true + + + false + + + false + + + + + + + + + + + + + + + + + + 0 + + + + + + 0 + 0 + + + + + 12 + true + + + + Config Selection + + + Qt::AlignmentFlag::AlignCenter + + + + + + 9 false + + + + + -1 + + + Common Config + + + + + + + + 10 + true + + - Use per-game configs + Common Config + + + Qt::AlignmentFlag::AlignCenter + + + true - + + + + + + 0 + 0 + + + + + 9 + false + + + + Use per-game configs + + + + + + + + + + 0 + + + + + + + L1 / LB + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + + + + + + L2 / LT + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + + + - - - 0 - + - - - - - L1 / LB - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - - - - L2 / LT - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - + + + Qt::Orientation::Vertical + + + QSizePolicy::Policy::Preferred + + + + 20 + 40 + + + - + + + 10 + - - - 10 - - - - - - true - - - - KBM Controls - - - - - - - true - - - - KBM Editor - - - - - - - - - - Back - - - - - - - - - - - - - - - - + - R1 / RB + Back - - - 5 - - - 5 - - - 5 - - - 5 - + - - - - - - - - - R2 / RT - - - - 5 - - - 5 - - - 5 - - - 5 - - - + @@ -788,62 +709,13 @@ - - - - 0 - 200 - - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 415 - 256 - - - - :/images/ps4_controller.png - - - true - - - Qt::AlignmentFlag::AlignBottom|Qt::AlignmentFlag::AlignHCenter - - - - - - - - - - 10 - - - QLayout::SizeConstraint::SetDefaultConstraint - + - + - L3 + R1 / RB - + 5 @@ -857,17 +729,17 @@ 5 - + - + - Options / Start + R2 / RT - + 5 @@ -881,31 +753,7 @@ 5 - - - - - - - - - R3 - - - - 5 - - - 5 - - - 5 - - - 5 - - - + @@ -915,22 +763,62 @@ - + + + + 0 + 200 + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 415 + 256 + + + + :/images/ps4_controller.png + + + true + + + Qt::AlignmentFlag::AlignBottom|Qt::AlignmentFlag::AlignHCenter + + + + + + + + - 5 + 10 + + + QLayout::SizeConstraint::SetDefaultConstraint - - - - 0 - 0 - - + - Face Buttons + L3 - + 5 @@ -944,244 +832,17 @@ 5 - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 0 - 0 - - - - - 124 - 0 - - - - - 0 - 16777215 - - - - Triangle / Y - - - - - - true - - - - 0 - 0 - - - - - - - - - - - - - - - - Square / X - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - - - - Circle / B - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - - - - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 124 - 0 - - - - - 124 - 16777215 - - - - Cross / A - - - - - - - - - - + - - - Qt::Orientation::Vertical - - - QSizePolicy::Policy::Maximum - - - - 20 - 40 - - - - - - - - - 0 - 0 - - + - Right Stick Deadzone (def:2, max:127) + Options / Start - - - - - - 0 - 0 - - - - Right Deadzone - - - Qt::AlignmentFlag::AlignRight|Qt::AlignmentFlag::AlignTrailing|Qt::AlignmentFlag::AlignVCenter - - - - - - - - 0 - 0 - - - - 1 - - - 127 - - - Qt::Orientation::Horizontal - - - - - - - - - - - 0 - 0 - - - - - 0 - 0 - - - - Right Stick - - + 5 @@ -1195,164 +856,637 @@ 5 - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 0 - 0 - - - - - 0 - 0 - - - - - 124 - 1231321 - - - - Up - - - - - - true - - - - - - - - + + + + + + + + R3 + + + + 5 + + + 5 + + + 5 + + + 5 + - - - - - Left - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - true - - - - - - - - - - Right - - - - 5 - - - 5 - - - 5 - - - 5 - - - - - - - - - - - - - - 0 - - - 0 - - - 0 - - - 0 - - - - - - 124 - 0 - - - - - 124 - 2121 - - - - Down - - - - - - true - - - - - - - - + + + + + + + + + + false + + + + Color Adjustment + + + + + + + + + false + + + + R: 000 + + + + + + + + 0 + 0 + + + + 255 + + + Qt::Orientation::Horizontal + + + + + + + + + + + + false + + + + G: 000 + + + + + + + + 0 + 0 + + + + 255 + + + Qt::Orientation::Horizontal + + + + + + + + + + + + false + + + + B: 255 + + + + + + + + 0 + 0 + + + + 255 + + + 255 + + + Qt::Orientation::Horizontal + + + + + + + + + + + + + + + + + false + + + + Override Lightbar Color + + + + + + + false + + + + Override Color + + + + + + + QFrame::Shape::StyledPanel + + + QFrame::Shadow::Raised + + + + + + + + + + + + + + + + 5 + + + + + + 0 + 0 + + + + Face Buttons + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 0 + 0 + + + + + 124 + 0 + + + + + 0 + 16777215 + + + + Triangle / Y + + + + + + true + + + + 0 + 0 + + + + + + + + + + + + + + + + Square / X + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + + + + + + Circle / B + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + + + + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 124 + 0 + + + + + 124 + 16777215 + + + + Cross / A + + + + + + + + + + + + + + + + + + Qt::Orientation::Vertical + + + QSizePolicy::Policy::Maximum + + + + 20 + 40 + + + + + + + + + 0 + 0 + + + + Right Stick Deadzone (def:2, max:127) + + + + + + + 0 + 0 + + + + Right Deadzone + + + Qt::AlignmentFlag::AlignRight|Qt::AlignmentFlag::AlignTrailing|Qt::AlignmentFlag::AlignVCenter + + + + + + + + 0 + 0 + + + + 1 + + + 127 + + + Qt::Orientation::Horizontal + + + + + + + + + + + 0 + 0 + + + + + 0 + 0 + + + + Right Stick + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 0 + 0 + + + + + 0 + 0 + + + + + 124 + 1231321 + + + + Up + + + + + + true + + + + + + + + + + + + + + + Left + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + true + + + + + + + + + + Right + + + + 5 + + + 5 + + + 5 + + + 5 + + + + + + + + + + + + + + 0 + + + 0 + + + 0 + + + 0 + + + + + + 124 + 0 + + + + + 124 + 2121 + + + + Down + + + + + + true + + + + + + + + + + + + diff --git a/src/qt_gui/game_grid_frame.cpp b/src/qt_gui/game_grid_frame.cpp index 6a42fb1d6..e06fea090 100644 --- a/src/qt_gui/game_grid_frame.cpp +++ b/src/qt_gui/game_grid_frame.cpp @@ -196,14 +196,28 @@ void GameGridFrame::SetGridBackgroundImage(int row, int column) { void GameGridFrame::RefreshGridBackgroundImage() { QPalette palette; if (!backgroundImage.isNull() && Config::getShowBackgroundImage()) { - palette.setBrush(QPalette::Base, - QBrush(backgroundImage.scaled(size(), Qt::IgnoreAspectRatio))); + QSize widgetSize = size(); + QPixmap scaledPixmap = + QPixmap::fromImage(backgroundImage) + .scaled(widgetSize, Qt::KeepAspectRatioByExpanding, Qt::SmoothTransformation); + int x = (widgetSize.width() - scaledPixmap.width()) / 2; + int y = (widgetSize.height() - scaledPixmap.height()) / 2; + QPixmap finalPixmap(widgetSize); + finalPixmap.fill(Qt::transparent); + QPainter painter(&finalPixmap); + painter.drawPixmap(x, y, scaledPixmap); + palette.setBrush(QPalette::Base, QBrush(finalPixmap)); } QColor transparentColor = QColor(135, 206, 235, 40); palette.setColor(QPalette::Highlight, transparentColor); this->setPalette(palette); } +void GameGridFrame::resizeEvent(QResizeEvent* event) { + QTableWidget::resizeEvent(event); + RefreshGridBackgroundImage(); +} + bool GameGridFrame::IsValidCellSelected() { return validCellSelected; } diff --git a/src/qt_gui/game_grid_frame.h b/src/qt_gui/game_grid_frame.h index 370b71dcb..14596f8e1 100644 --- a/src/qt_gui/game_grid_frame.h +++ b/src/qt_gui/game_grid_frame.h @@ -3,6 +3,7 @@ #pragma once +#include #include #include "background_music_player.h" @@ -21,6 +22,7 @@ Q_SIGNALS: public Q_SLOTS: void SetGridBackgroundImage(int row, int column); void RefreshGridBackgroundImage(); + void resizeEvent(QResizeEvent* event); void PlayBackgroundMusic(QString path); void onCurrentCellChanged(int currentRow, int currentColumn, int previousRow, int previousColumn); diff --git a/src/qt_gui/game_install_dialog.h b/src/qt_gui/game_install_dialog.h index 0a4e29357..938f0e1f3 100644 --- a/src/qt_gui/game_install_dialog.h +++ b/src/qt_gui/game_install_dialog.h @@ -11,6 +11,7 @@ class QLineEdit; class GameInstallDialog final : public QDialog { + Q_OBJECT public: GameInstallDialog(); ~GameInstallDialog(); diff --git a/src/qt_gui/game_list_frame.cpp b/src/qt_gui/game_list_frame.cpp index 2caae35b0..4c0607571 100644 --- a/src/qt_gui/game_list_frame.cpp +++ b/src/qt_gui/game_list_frame.cpp @@ -200,14 +200,28 @@ void GameListFrame::SetListBackgroundImage(QTableWidgetItem* item) { void GameListFrame::RefreshListBackgroundImage() { QPalette palette; if (!backgroundImage.isNull() && Config::getShowBackgroundImage()) { - palette.setBrush(QPalette::Base, - QBrush(backgroundImage.scaled(size(), Qt::IgnoreAspectRatio))); + QSize widgetSize = size(); + QPixmap scaledPixmap = + QPixmap::fromImage(backgroundImage) + .scaled(widgetSize, Qt::KeepAspectRatioByExpanding, Qt::SmoothTransformation); + int x = (widgetSize.width() - scaledPixmap.width()) / 2; + int y = (widgetSize.height() - scaledPixmap.height()) / 2; + QPixmap finalPixmap(widgetSize); + finalPixmap.fill(Qt::transparent); + QPainter painter(&finalPixmap); + painter.drawPixmap(x, y, scaledPixmap); + palette.setBrush(QPalette::Base, QBrush(finalPixmap)); } QColor transparentColor = QColor(135, 206, 235, 40); palette.setColor(QPalette::Highlight, transparentColor); this->setPalette(palette); } +void GameListFrame::resizeEvent(QResizeEvent* event) { + QTableWidget::resizeEvent(event); + RefreshListBackgroundImage(); +} + void GameListFrame::SortNameAscending(int columnIndex) { std::sort(m_game_info->m_games.begin(), m_game_info->m_games.end(), [columnIndex](const GameInfo& a, const GameInfo& b) { diff --git a/src/qt_gui/game_list_frame.h b/src/qt_gui/game_list_frame.h index b2e5f1e2f..782db6bae 100644 --- a/src/qt_gui/game_list_frame.h +++ b/src/qt_gui/game_list_frame.h @@ -30,6 +30,7 @@ Q_SIGNALS: public Q_SLOTS: void SetListBackgroundImage(QTableWidgetItem* item); void RefreshListBackgroundImage(); + void resizeEvent(QResizeEvent* event); void SortNameAscending(int columnIndex); void SortNameDescending(int columnIndex); void PlayBackgroundMusic(QTableWidgetItem* item); diff --git a/src/qt_gui/settings_dialog.cpp b/src/qt_gui/settings_dialog.cpp index 1598b0640..bebb16c9a 100644 --- a/src/qt_gui/settings_dialog.cpp +++ b/src/qt_gui/settings_dialog.cpp @@ -536,7 +536,7 @@ void SettingsDialog::updateNoteTextEdit(const QString& elementName) { } else if (elementName == "fullscreenCheckBox") { text = tr("Enable Full Screen:\\nAutomatically puts the game window into full-screen mode.\\nThis can be toggled by pressing the F11 key."); } else if (elementName == "separateUpdatesCheckBox") { - text = tr("Enable Separate Update Folder:\\nEnables installing game updates into a separate folder for easy management.\\nThis can be manually created by adding the extracted update to the game folder with the name 'CUSA00000-UPDATE' where the CUSA ID matches the game's ID."); + text = tr("Enable Separate Update Folder:\\nEnables installing game updates into a separate folder for easy management.\\nThis can be manually created by adding the extracted update to the game folder with the name \"CUSA00000-UPDATE\" where the CUSA ID matches the game's ID."); } else if (elementName == "showSplashCheckBox") { text = tr("Show Splash Screen:\\nShows the game's splash screen (a special image) while the game is starting."); } else if (elementName == "discordRPCCheckbox") { @@ -548,8 +548,8 @@ void SettingsDialog::updateNoteTextEdit(const QString& elementName) { } else if (elementName == "logTypeGroupBox") { text = tr("Log Type:\\nSets whether to synchronize the output of the log window for performance. May have adverse effects on emulation."); } else if (elementName == "logFilter") { - text = tr("Log Filter:\nFilters the log to only print specific information.\nExamples: 'Core:Trace' 'Lib.Pad:Debug Common.Filesystem:Error' '*:Critical'\\nLevels: Trace, Debug, Info, Warning, Error, Critical - in this order, a specific level silences all levels preceding it in the list and logs every level after it."); -#ifdef ENABLE_UPDATER + text = tr("Log Filter:\\nFilters the log to only print specific information.\\nExamples: \"Core:Trace\" \"Lib.Pad:Debug Common.Filesystem:Error\" \"*:Critical\"\\nLevels: Trace, Debug, Info, Warning, Error, Critical - in this order, a specific level silences all levels preceding it in the list and logs every level after it."); + #ifdef ENABLE_UPDATER } else if (elementName == "updaterGroupBox") { text = tr("Update:\\nRelease: Official versions released every month that may be very outdated, but are more reliable and tested.\\nNightly: Development versions that have all the latest features and fixes, but may contain bugs and are less stable."); #endif @@ -562,7 +562,7 @@ void SettingsDialog::updateNoteTextEdit(const QString& elementName) { } else if (elementName == "disableTrophycheckBox") { text = tr("Disable Trophy Pop-ups:\\nDisable in-game trophy notifications. Trophy progress can still be tracked using the Trophy Viewer (right-click the game in the main window)."); } else if (elementName == "enableCompatibilityCheckBox") { - text = tr("Display Compatibility Data:\\nDisplays game compatibility information in table view. Enable 'Update Compatibility On Startup' to get up-to-date information."); + text = tr("Display Compatibility Data:\\nDisplays game compatibility information in table view. Enable \"Update Compatibility On Startup\" to get up-to-date information."); } else if (elementName == "checkCompatibilityOnStartupCheckBox") { text = tr("Update Compatibility On Startup:\\nAutomatically update the compatibility database when shadPS4 starts."); } else if (elementName == "updateCompatibilityButton") { @@ -580,7 +580,7 @@ void SettingsDialog::updateNoteTextEdit(const QString& elementName) { // Graphics if (elementName == "graphicsAdapterGroupBox") { - text = tr("Graphics Device:\\nOn multiple GPU systems, select the GPU the emulator will use from the drop down list,\\nor select 'Auto Select' to automatically determine it."); + text = tr("Graphics Device:\\nOn multiple GPU systems, select the GPU the emulator will use from the drop down list,\\nor select \"Auto Select\" to automatically determine it."); } else if (elementName == "widthGroupBox" || elementName == "heightGroupBox") { text = tr("Width/Height:\\nSets the size of the emulator window at launch, which can be resized during gameplay.\\nThis is different from the in-game resolution."); } else if (elementName == "heightDivider") { diff --git a/src/qt_gui/translations/ja_JP.ts b/src/qt_gui/translations/ja_JP.ts index fd1de8f78..502070bb5 100644 --- a/src/qt_gui/translations/ja_JP.ts +++ b/src/qt_gui/translations/ja_JP.ts @@ -11,7 +11,7 @@ shadPS4 - shadPS4 + shadPS4 shadPS4 is an experimental open-source emulator for the PlayStation 4. @@ -411,35 +411,35 @@ ControlSettings Configure Controls - Configure Controls + コントロール設定 Control Settings - Control Settings + 操作設定 D-Pad - D-Pad + 十字キー Up - Up + 上 Left - Left + 左 Right - Right + 右 Down - Down + 下 Left Stick Deadzone (def:2 max:127) - Left Stick Deadzone (def:2 max:127) + 左スティックデッドゾーン（既定:2 最大:127） Left Deadzone @@ -447,7 +447,7 @@ Left Stick - Left Stick + 左スティック Config Selection @@ -463,11 +463,11 @@ L1 / LB - L1 / LB + L1 / LB L2 / LT - L2 / LT + L2 / LT KBM Controls @@ -483,23 +483,23 @@ R1 / RB - R1 / RB + R1 / RB R2 / RT - R2 / RT + R2 / RT L3 - L3 + L3 Options / Start - Options / Start + Options / Start R3 - R3 + R3 Face Buttons @@ -507,23 +507,23 @@ Triangle / Y - Triangle / Y + 三角 / Y Square / X - Square / X + 四角 / X Circle / B - Circle / B + 丸 / B Cross / A - Cross / A + バツ / A Right Stick Deadzone (def:2, max:127) - Right Stick Deadzone (def:2, max:127) + 右スティックデッドゾーン（既定:2, 最大:127） Right Deadzone @@ -531,7 +531,7 @@ Right Stick - Right Stick + 右スティック @@ -576,7 +576,7 @@ Directory to install DLC - Directory to install DLC + DLCをインストールするディレクトリ @@ -595,7 +595,7 @@ Compatibility - Compatibility + 互換性 Region @@ -674,23 +674,23 @@ GameListUtils B - B + B KB - KB + KB MB - MB + MB GB - GB + GB TB - TB + TB @@ -741,11 +741,11 @@ Copy Version - Copy Version + バージョンをコピー Copy Size - Copy Size + サイズをコピー Copy All @@ -821,7 +821,7 @@ DLC - DLC + DLC Delete %1 @@ -833,23 +833,23 @@ Open Update Folder - Open Update Folder + アップデートフォルダを開く Delete Save Data - Delete Save Data + セーブデータを削除 This game has no update folder to open! - This game has no update folder to open! + このゲームにはアップデートフォルダがありません！ Failed to convert icon. - Failed to convert icon. + アイコンの変換に失敗しました。 This game has no save data to delete! - This game has no save data to delete! + このゲームには削除するセーブデータがありません！ Save Data @@ -872,7 +872,7 @@ Delete PKG File on Install - Delete PKG File on Install + インストール時にPKGファイルを削除 @@ -1151,15 +1151,15 @@ Run Game - Run Game + ゲームを実行 Eboot.bin file not found - Eboot.bin file not found + Eboot.bin ファイルが見つかりません PKG File (*.PKG *.pkg) - PKG File (*.PKG *.pkg) + PKGファイル (*.PKG *.pkg) PKG is a patch or DLC, please install the game first! @@ -1167,11 +1167,11 @@ Game is already running! - Game is already running! + ゲームは既に実行されています！ shadPS4 - shadPS4 + shadPS4 @@ -1238,7 +1238,7 @@ Package - Package + パッケージ @@ -1393,7 +1393,7 @@ Enable HDR - Enable HDR + HDRを有効化 Paths @@ -1493,7 +1493,7 @@ Opacity - Opacity + 透明度 Play title music @@ -1729,7 +1729,7 @@ Borderless - Borderless + ボーダーレス True @@ -1737,19 +1737,19 @@ Release - Release + Release Nightly - Nightly + Nightly Set the volume of the background music. - Set the volume of the background music. + バックグラウンドミュージックの音量を設定します。 Enable Motion Controls - Enable Motion Controls + モーションコントロールを有効にする Save Data Path @@ -1761,11 +1761,11 @@ async - async + 非同期 sync - sync + 同期 Auto Select diff --git a/src/qt_gui/translations/no_NO.ts b/src/qt_gui/translations/nb_NO.ts similarity index 99% rename from src/qt_gui/translations/no_NO.ts rename to src/qt_gui/translations/nb_NO.ts index 2613f63b0..934612683 100644 --- a/src/qt_gui/translations/no_NO.ts +++ b/src/qt_gui/translations/nb_NO.ts @@ -2,7 +2,7 @@ - + AboutDialog diff --git a/src/qt_gui/translations/pt_BR.ts b/src/qt_gui/translations/pt_BR.ts index 33f76764f..be9c4d11b 100644 --- a/src/qt_gui/translations/pt_BR.ts +++ b/src/qt_gui/translations/pt_BR.ts @@ -1309,7 +1309,7 @@ Logger - Registro-Log + Registros de Log Log Type @@ -1497,7 +1497,7 @@ Play title music - Reproduzir música de abertura + Reproduzir Música do Título Update Compatibility Database On Startup @@ -1573,7 +1573,7 @@ Log Type:\nSets whether to synchronize the output of the log window for performance. May have adverse effects on emulation. - Tipo do Registro:\nDetermina se a saída da janela de log deve ser sincronizada por motivos de desempenho. Pode impactar negativamente a emulação. + Tipo de Registro:\nDetermina se a saída da janela de log deve ser sincronizada por motivos de desempenho. Pode impactar negativamente a emulação. Log Filter:\nFilters the log to only print specific information.\nExamples: "Core:Trace" "Lib.Pad:Debug Common.Filesystem:Error" "*:Critical"\nLevels: Trace, Debug, Info, Warning, Error, Critical - in this order, a specific level silences all levels preceding it in the list and logs every level after it. @@ -1585,7 +1585,7 @@ Background Image:\nControl the opacity of the game background image. - Imagem de fundo:\nControle a opacidade da imagem de fundo do jogo. + Imagem de Fundo:\nControla a opacidade da imagem de fundo do jogo. Play Title Music:\nIf a game supports it, enable playing special music when selecting the game in the GUI. @@ -1705,7 +1705,7 @@ Crash Diagnostics:\nCreates a .yaml file with info about the Vulkan state at the time of crashing.\nUseful for debugging 'Device lost' errors. If you have this enabled, you should enable Host AND Guest Debug Markers.\nDoes not work on Intel GPUs.\nYou need Vulkan Validation Layers enabled and the Vulkan SDK for this to work. - Diagnósticos de Falha:\nCria um arquivo .yaml com informações sobre o estado do Vulkan no momento da falha.\nÚtil para depuração de erros de 'Device lost'. Se você tiver isto habilitado, você deve habilitar os Marcadores de Depuração de Host E DE Convidado.\nNão funciona em GPUs Intel.\nVocê precisa ter as Camadas de Validação Vulkan habilitadas e o Vulkan SDK para que isso funcione. + Diagnóstico de Falhas:\nCria um arquivo .yaml com informações sobre o estado do Vulkan no momento da falha.\nÚtil para depuração de erros de 'Device lost'. Se isto estiver ativado, você deve habilitar os Marcadores de Depuração de Host E DE Convidado.\nNão funciona em GPUs Intel.\nVocê precisa ter as Camadas de Validação Vulkan habilitadas e o Vulkan SDK para que isso funcione. Copy GPU Buffers:\nGets around race conditions involving GPU submits.\nMay or may not help with PM4 type 0 crashes. @@ -1721,11 +1721,11 @@ Save Data Path:\nThe folder where game save data will be saved. - Diretório dos Dados Salvos:\nA pasta que onde os dados de salvamento de jogo serão salvos. + Caminho dos Dados Salvos:\nA pasta que onde os dados de salvamento de jogo serão salvos. Browse:\nBrowse for a folder to set as the save data path. - Navegar:\nProcure uma pasta para definir como o caminho para salvar dados. + Procurar:\nProcure uma pasta para definir como o caminho para salvar dados. Borderless diff --git a/src/qt_gui/translations/pt_PT.ts b/src/qt_gui/translations/pt_PT.ts new file mode 100644 index 000000000..1a63c88fd --- /dev/null +++ b/src/qt_gui/translations/pt_PT.ts @@ -0,0 +1,1790 @@ + + + + + + AboutDialog + + About shadPS4 + Sobre o shadPS4 + + + shadPS4 + shadPS4 + + + shadPS4 is an experimental open-source emulator for the PlayStation 4. + shadPS4 é um emulador de código aberto experimental para o PlayStation 4. + + + This software should not be used to play games you have not legally obtained. + Este programa não deve ser usado para jogar títulos não obtidos legalmente. + + + + CheatsPatches + + Cheats / Patches for + Cheats / Patches for + + + Cheats/Patches are experimental.\nUse with caution.\n\nDownload cheats individually by selecting the repository and clicking the download button.\nIn the Patches tab, you can download all patches at once, choose which ones you want to use, and save your selection.\n\nSince we do not develop the Cheats/Patches,\nplease report issues to the cheat author.\n\nCreated a new cheat? Visit:\n + Cheats/Patches are experimental.\nUse with caution.\n\nDownload cheats individually by selecting the repository and clicking the download button.\nIn the Patches tab, you can download all patches at once, choose which ones you want to use, and save your selection.\n\nSince we do not develop the Cheats/Patches,\nplease report issues to the cheat author.\n\nCreated a new cheat? Visit:\n + + + No Image Available + No Image Available + + + Serial: + Serial: + + + Version: + Version: + + + Size: + Size: + + + Select Cheat File: + Select Cheat File: + + + Repository: + Repository: + + + Download Cheats + Download Cheats + + + Delete File + Delete File + + + No files selected. + No files selected. + + + You can delete the cheats you don't want after downloading them. + You can delete the cheats you don't want after downloading them. + + + Do you want to delete the selected file?\n%1 + Do you want to delete the selected file?\n%1 + + + Select Patch File: + Select Patch File: + + + Download Patches + Download Patches + + + Save + Save + + + Cheats + Cheats + + + Patches + Patches + + + Error + Error + + + No patch selected. + No patch selected. + + + Unable to open files.json for reading. + Unable to open files.json for reading. + + + No patch file found for the current serial. + No patch file found for the current serial. + + + Unable to open the file for reading. + Unable to open the file for reading. + + + Unable to open the file for writing. + Unable to open the file for writing. + + + Failed to parse XML: + Failed to parse XML: + + + Success + Success + + + Options saved successfully. + Options saved successfully. + + + Invalid Source + Invalid Source + + + The selected source is invalid. + The selected source is invalid. + + + File Exists + File Exists + + + File already exists. Do you want to replace it? + File already exists. Do you want to replace it? + + + Failed to save file: + Failed to save file: + + + Failed to download file: + Failed to download file: + + + Cheats Not Found + Cheats Not Found + + + No Cheats found for this game in this version of the selected repository,try another repository or a different version of the game. + No Cheats found for this game in this version of the selected repository,try another repository or a different version of the game. + + + Cheats Downloaded Successfully + Cheats Downloaded Successfully + + + You have successfully downloaded the cheats for this version of the game from the selected repository. You can try downloading from another repository, if it is available it will also be possible to use it by selecting the file from the list. + You have successfully downloaded the cheats for this version of the game from the selected repository. You can try downloading from another repository, if it is available it will also be possible to use it by selecting the file from the list. + + + Failed to save: + Failed to save: + + + Failed to download: + Failed to download: + + + Download Complete + Download Complete + + + Patches Downloaded Successfully! All Patches available for all games have been downloaded, there is no need to download them individually for each game as happens in Cheats. If the patch does not appear, it may be that it does not exist for the specific serial and version of the game. + Patches Downloaded Successfully! All Patches available for all games have been downloaded, there is no need to download them individually for each game as happens in Cheats. If the patch does not appear, it may be that it does not exist for the specific serial and version of the game. + + + Failed to parse JSON data from HTML. + Failed to parse JSON data from HTML. + + + Failed to retrieve HTML page. + Failed to retrieve HTML page. + + + The game is in version: %1 + The game is in version: %1 + + + The downloaded patch only works on version: %1 + The downloaded patch only works on version: %1 + + + You may need to update your game. + You may need to update your game. + + + Incompatibility Notice + Incompatibility Notice + + + Failed to open file: + Failed to open file: + + + XML ERROR: + XML ERROR: + + + Failed to open files.json for writing + Failed to open files.json for writing + + + Author: + Author: + + + Directory does not exist: + Directory does not exist: + + + Failed to open files.json for reading. + Failed to open files.json for reading. + + + Name: + Name: + + + Can't apply cheats before the game is started + Can't apply cheats before the game is started + + + Close + Close + + + + CheckUpdate + + Auto Updater + Auto Updater + + + Error + Error + + + Network error: + Network error: + + + The Auto Updater allows up to 60 update checks per hour.\nYou have reached this limit. Please try again later. + The Auto Updater allows up to 60 update checks per hour.\nYou have reached this limit. Please try again later. + + + Failed to parse update information. + Failed to parse update information. + + + No pre-releases found. + No pre-releases found. + + + Invalid release data. + Invalid release data. + + + No download URL found for the specified asset. + No download URL found for the specified asset. + + + Your version is already up to date! + Your version is already up to date! + + + Update Available + Update Available + + + Update Channel + Update Channel + + + Current Version + Current Version + + + Latest Version + Latest Version + + + Do you want to update? + Do you want to update? + + + Show Changelog + Show Changelog + + + Check for Updates at Startup + Check for Updates at Startup + + + Update + Update + + + No + No + + + Hide Changelog + Hide Changelog + + + Changes + Changes + + + Network error occurred while trying to access the URL + Network error occurred while trying to access the URL + + + Download Complete + Download Complete + + + The update has been downloaded, press OK to install. + The update has been downloaded, press OK to install. + + + Failed to save the update file at + Failed to save the update file at + + + Starting Update... + Starting Update... + + + Failed to create the update script file + Failed to create the update script file + + + + CompatibilityInfoClass + + Fetching compatibility data, please wait + Fetching compatibility data, please wait + + + Cancel + Cancel + + + Loading... + Loading... + + + Error + Error + + + Unable to update compatibility data! Try again later. + Unable to update compatibility data! Try again later. + + + Unable to open compatibility_data.json for writing. + Unable to open compatibility_data.json for writing. + + + Unknown + Unknown + + + Nothing + Nothing + + + Boots + Boots + + + Menus + Menus + + + Ingame + Ingame + + + Playable + Playable + + + + ControlSettings + + Configure Controls + Configure Controls + + + Control Settings + Control Settings + + + D-Pad + D-Pad + + + Up + Up + + + Left + Left + + + Right + Right + + + Down + Down + + + Left Stick Deadzone (def:2 max:127) + Left Stick Deadzone (def:2 max:127) + + + Left Deadzone + Left Deadzone + + + Left Stick + Left Stick + + + Config Selection + Config Selection + + + Common Config + Common Config + + + Use per-game configs + Use per-game configs + + + L1 / LB + L1 / LB + + + L2 / LT + L2 / LT + + + KBM Controls + KBM Controls + + + KBM Editor + KBM Editor + + + Back + Back + + + R1 / RB + R1 / RB + + + R2 / RT + R2 / RT + + + L3 + L3 + + + Options / Start + Options / Start + + + R3 + R3 + + + Face Buttons + Face Buttons + + + Triangle / Y + Triangle / Y + + + Square / X + Square / X + + + Circle / B + Circle / B + + + Cross / A + Cross / A + + + Right Stick Deadzone (def:2, max:127) + Right Stick Deadzone (def:2, max:127) + + + Right Deadzone + Right Deadzone + + + Right Stick + Right Stick + + + + ElfViewer + + Open Folder + Open Folder + + + + GameInfoClass + + Loading game list, please wait :3 + Loading game list, please wait :3 + + + Cancel + Cancel + + + Loading... + Loading... + + + + GameInstallDialog + + shadPS4 - Choose directory + shadPS4 - Choose directory + + + Directory to install games + Directory to install games + + + Browse + Browse + + + Error + Error + + + Directory to install DLC + Directory to install DLC + + + + GameListFrame + + Icon + Icon + + + Name + Name + + + Serial + Serial + + + Compatibility + Compatibility + + + Region + Region + + + Firmware + Firmware + + + Size + Size + + + Version + Version + + + Path + Path + + + Play Time + Play Time + + + Never Played + Never Played + + + h + h + + + m + m + + + s + s + + + Compatibility is untested + Compatibility is untested + + + Game does not initialize properly / crashes the emulator + Game does not initialize properly / crashes the emulator + + + Game boots, but only displays a blank screen + Game boots, but only displays a blank screen + + + Game displays an image but does not go past the menu + Game displays an image but does not go past the menu + + + Game has game-breaking glitches or unplayable performance + Game has game-breaking glitches or unplayable performance + + + Game can be completed with playable performance and no major glitches + Game can be completed with playable performance and no major glitches + + + Click to see details on github + Click to see details on github + + + Last updated + Last updated + + + + GameListUtils + + B + B + + + KB + KB + + + MB + MB + + + GB + GB + + + TB + TB + + + + GuiContextMenus + + Create Shortcut + Create Shortcut + + + Cheats / Patches + Cheats / Patches + + + SFO Viewer + SFO Viewer + + + Trophy Viewer + Trophy Viewer + + + Open Folder... + Open Folder... + + + Open Game Folder + Open Game Folder + + + Open Save Data Folder + Open Save Data Folder + + + Open Log Folder + Open Log Folder + + + Copy info... + Copy info... + + + Copy Name + Copy Name + + + Copy Serial + Copy Serial + + + Copy Version + Copy Version + + + Copy Size + Copy Size + + + Copy All + Copy All + + + Delete... + Delete... + + + Delete Game + Delete Game + + + Delete Update + Delete Update + + + Delete DLC + Delete DLC + + + Compatibility... + Compatibility... + + + Update database + Update database + + + View report + View report + + + Submit a report + Submit a report + + + Shortcut creation + Shortcut creation + + + Shortcut created successfully! + Shortcut created successfully! + + + Error + Error + + + Error creating shortcut! + Error creating shortcut! + + + Install PKG + Install PKG + + + Game + Game + + + This game has no update to delete! + This game has no update to delete! + + + Update + Update + + + This game has no DLC to delete! + This game has no DLC to delete! + + + DLC + DLC + + + Delete %1 + Delete %1 + + + Are you sure you want to delete %1's %2 directory? + Are you sure you want to delete %1's %2 directory? + + + Open Update Folder + Open Update Folder + + + Delete Save Data + Delete Save Data + + + This game has no update folder to open! + This game has no update folder to open! + + + Failed to convert icon. + Failed to convert icon. + + + This game has no save data to delete! + This game has no save data to delete! + + + Save Data + Save Data + + + + InstallDirSelect + + shadPS4 - Choose directory + shadPS4 - Choose directory + + + Select which directory you want to install to. + Select which directory you want to install to. + + + Install All Queued to Selected Folder + Install All Queued to Selected Folder + + + Delete PKG File on Install + Delete PKG File on Install + + + + MainWindow + + Open/Add Elf Folder + Open/Add Elf Folder + + + Install Packages (PKG) + Install Packages (PKG) + + + Boot Game + Boot Game + + + Check for Updates + Check for Updates + + + About shadPS4 + About shadPS4 + + + Configure... + Configure... + + + Install application from a .pkg file + Install application from a .pkg file + + + Recent Games + Recent Games + + + Open shadPS4 Folder + Open shadPS4 Folder + + + Exit + Exit + + + Exit shadPS4 + Exit shadPS4 + + + Exit the application. + Exit the application. + + + Show Game List + Show Game List + + + Game List Refresh + Game List Refresh + + + Tiny + Tiny + + + Small + Small + + + Medium + Medium + + + Large + Large + + + List View + List View + + + Grid View + Grid View + + + Elf Viewer + Elf Viewer + + + Game Install Directory + Game Install Directory + + + Download Cheats/Patches + Download Cheats/Patches + + + Dump Game List + Dump Game List + + + PKG Viewer + PKG Viewer + + + Search... + Search... + + + File + File + + + View + View + + + Game List Icons + Game List Icons + + + Game List Mode + Game List Mode + + + Settings + Settings + + + Utils + Utils + + + Themes + Themes + + + Help + Help + + + Dark + Dark + + + Light + Light + + + Green + Green + + + Blue + Blue + + + Violet + Violet + + + toolBar + toolBar + + + Game List + Game List + + + * Unsupported Vulkan Version + * Unsupported Vulkan Version + + + Download Cheats For All Installed Games + Download Cheats For All Installed Games + + + Download Patches For All Games + Download Patches For All Games + + + Download Complete + Download Complete + + + You have downloaded cheats for all the games you have installed. + You have downloaded cheats for all the games you have installed. + + + Patches Downloaded Successfully! + Patches Downloaded Successfully! + + + All Patches available for all games have been downloaded. + All Patches available for all games have been downloaded. + + + Games: + Games: + + + ELF files (*.bin *.elf *.oelf) + ELF files (*.bin *.elf *.oelf) + + + Game Boot + Game Boot + + + Only one file can be selected! + Only one file can be selected! + + + PKG Extraction + PKG Extraction + + + Patch detected! + Patch detected! + + + PKG and Game versions match: + PKG and Game versions match: + + + Would you like to overwrite? + Would you like to overwrite? + + + PKG Version %1 is older than installed version: + PKG Version %1 is older than installed version: + + + Game is installed: + Game is installed: + + + Would you like to install Patch: + Would you like to install Patch: + + + DLC Installation + DLC Installation + + + Would you like to install DLC: %1? + Would you like to install DLC: %1? + + + DLC already installed: + DLC already installed: + + + Game already installed + Game already installed + + + PKG ERROR + PKG ERROR + + + Extracting PKG %1/%2 + Extracting PKG %1/%2 + + + Extraction Finished + Extraction Finished + + + Game successfully installed at %1 + Game successfully installed at %1 + + + File doesn't appear to be a valid PKG file + File doesn't appear to be a valid PKG file + + + Run Game + Run Game + + + Eboot.bin file not found + Eboot.bin file not found + + + PKG File (*.PKG *.pkg) + PKG File (*.PKG *.pkg) + + + PKG is a patch or DLC, please install the game first! + PKG is a patch or DLC, please install the game first! + + + Game is already running! + Game is already running! + + + shadPS4 + shadPS4 + + + + PKGViewer + + Open Folder + Open Folder + + + PKG ERROR + PKG ERROR + + + Name + Name + + + Serial + Serial + + + Installed + Installed + + + Size + Size + + + Category + Category + + + Type + Type + + + App Ver + App Ver + + + FW + FW + + + Region + Region + + + Flags + Flags + + + Path + Path + + + File + File + + + Unknown + Unknown + + + Package + Package + + + + SettingsDialog + + Settings + Settings + + + General + General + + + System + System + + + Console Language + Console Language + + + Emulator Language + Emulator Language + + + Emulator + Emulator + + + Enable Fullscreen + Enable Fullscreen + + + Fullscreen Mode + Fullscreen Mode + + + Enable Separate Update Folder + Enable Separate Update Folder + + + Default tab when opening settings + Default tab when opening settings + + + Show Game Size In List + Show Game Size In List + + + Show Splash + Show Splash + + + Enable Discord Rich Presence + Enable Discord Rich Presence + + + Username + Username + + + Trophy Key + Trophy Key + + + Trophy + Trophy + + + Logger + Logger + + + Log Type + Log Type + + + Log Filter + Log Filter + + + Open Log Location + Open Log Location + + + Input + Input + + + Cursor + Cursor + + + Hide Cursor + Hide Cursor + + + Hide Cursor Idle Timeout + Hide Cursor Idle Timeout + + + s + s + + + Controller + Controller + + + Back Button Behavior + Back Button Behavior + + + Graphics + Graphics + + + GUI + GUI + + + User + User + + + Graphics Device + Graphics Device + + + Width + Width + + + Height + Height + + + Vblank Divider + Vblank Divider + + + Advanced + Advanced + + + Enable Shaders Dumping + Enable Shaders Dumping + + + Enable NULL GPU + Enable NULL GPU + + + Enable HDR + Enable HDR + + + Paths + Paths + + + Game Folders + Game Folders + + + Add... + Add... + + + Remove + Remove + + + Debug + Debug + + + Enable Debug Dumping + Enable Debug Dumping + + + Enable Vulkan Validation Layers + Enable Vulkan Validation Layers + + + Enable Vulkan Synchronization Validation + Enable Vulkan Synchronization Validation + + + Enable RenderDoc Debugging + Enable RenderDoc Debugging + + + Enable Crash Diagnostics + Enable Crash Diagnostics + + + Collect Shaders + Collect Shaders + + + Copy GPU Buffers + Copy GPU Buffers + + + Host Debug Markers + Host Debug Markers + + + Guest Debug Markers + Guest Debug Markers + + + Update + Update + + + Check for Updates at Startup + Check for Updates at Startup + + + Always Show Changelog + Always Show Changelog + + + Update Channel + Update Channel + + + Check for Updates + Check for Updates + + + GUI Settings + GUI Settings + + + Title Music + Title Music + + + Disable Trophy Pop-ups + Disable Trophy Pop-ups + + + Background Image + Background Image + + + Show Background Image + Show Background Image + + + Opacity + Opacity + + + Play title music + Play title music + + + Update Compatibility Database On Startup + Update Compatibility Database On Startup + + + Game Compatibility + Game Compatibility + + + Display Compatibility Data + Display Compatibility Data + + + Update Compatibility Database + Update Compatibility Database + + + Volume + Volume + + + Save + Save + + + Apply + Apply + + + Restore Defaults + Restore Defaults + + + Close + Close + + + Point your mouse at an option to display its description. + Point your mouse at an option to display its description. + + + Console Language:\nSets the language that the PS4 game uses.\nIt's recommended to set this to a language the game supports, which will vary by region. + Console Language:\nSets the language that the PS4 game uses.\nIt's recommended to set this to a language the game supports, which will vary by region. + + + Emulator Language:\nSets the language of the emulator's user interface. + Emulator Language:\nSets the language of the emulator's user interface. + + + Enable Full Screen:\nAutomatically puts the game window into full-screen mode.\nThis can be toggled by pressing the F11 key. + Enable Full Screen:\nAutomatically puts the game window into full-screen mode.\nThis can be toggled by pressing the F11 key. + + + Enable Separate Update Folder:\nEnables installing game updates into a separate folder for easy management.\nThis can be manually created by adding the extracted update to the game folder with the name "CUSA00000-UPDATE" where the CUSA ID matches the game's ID. + Enable Separate Update Folder:\nEnables installing game updates into a separate folder for easy management.\nThis can be manually created by adding the extracted update to the game folder with the name "CUSA00000-UPDATE" where the CUSA ID matches the game's ID. + + + Show Splash Screen:\nShows the game's splash screen (a special image) while the game is starting. + Show Splash Screen:\nShows the game's splash screen (a special image) while the game is starting. + + + Enable Discord Rich Presence:\nDisplays the emulator icon and relevant information on your Discord profile. + Enable Discord Rich Presence:\nDisplays the emulator icon and relevant information on your Discord profile. + + + Username:\nSets the PS4's account username, which may be displayed by some games. + Username:\nSets the PS4's account username, which may be displayed by some games. + + + Trophy Key:\nKey used to decrypt trophies. Must be obtained from your jailbroken console.\nMust contain only hex characters. + Trophy Key:\nKey used to decrypt trophies. Must be obtained from your jailbroken console.\nMust contain only hex characters. + + + Log Type:\nSets whether to synchronize the output of the log window for performance. May have adverse effects on emulation. + Log Type:\nSets whether to synchronize the output of the log window for performance. May have adverse effects on emulation. + + + Log Filter:\nFilters the log to only print specific information.\nExamples: "Core:Trace" "Lib.Pad:Debug Common.Filesystem:Error" "*:Critical"\nLevels: Trace, Debug, Info, Warning, Error, Critical - in this order, a specific level silences all levels preceding it in the list and logs every level after it. + Log Filter:\nFilters the log to only print specific information.\nExamples: "Core:Trace" "Lib.Pad:Debug Common.Filesystem:Error" "*:Critical"\nLevels: Trace, Debug, Info, Warning, Error, Critical - in this order, a specific level silences all levels preceding it in the list and logs every level after it. + + + Update:\nRelease: Official versions released every month that may be very outdated, but are more reliable and tested.\nNightly: Development versions that have all the latest features and fixes, but may contain bugs and are less stable. + Update:\nRelease: Official versions released every month that may be very outdated, but are more reliable and tested.\nNightly: Development versions that have all the latest features and fixes, but may contain bugs and are less stable. + + + Background Image:\nControl the opacity of the game background image. + Background Image:\nControl the opacity of the game background image. + + + Play Title Music:\nIf a game supports it, enable playing special music when selecting the game in the GUI. + Play Title Music:\nIf a game supports it, enable playing special music when selecting the game in the GUI. + + + Disable Trophy Pop-ups:\nDisable in-game trophy notifications. Trophy progress can still be tracked using the Trophy Viewer (right-click the game in the main window). + Disable Trophy Pop-ups:\nDisable in-game trophy notifications. Trophy progress can still be tracked using the Trophy Viewer (right-click the game in the main window). + + + Hide Cursor:\nChoose when the cursor will disappear:\nNever: You will always see the mouse.\nidle: Set a time for it to disappear after being idle.\nAlways: you will never see the mouse. + Hide Cursor:\nChoose when the cursor will disappear:\nNever: You will always see the mouse.\nidle: Set a time for it to disappear after being idle.\nAlways: you will never see the mouse. + + + Hide Idle Cursor Timeout:\nThe duration (seconds) after which the cursor that has been idle hides itself. + Hide Idle Cursor Timeout:\nThe duration (seconds) after which the cursor that has been idle hides itself. + + + Back Button Behavior:\nSets the controller's back button to emulate tapping the specified position on the PS4 touchpad. + Back Button Behavior:\nSets the controller's back button to emulate tapping the specified position on the PS4 touchpad. + + + Display Compatibility Data:\nDisplays game compatibility information in table view. Enable "Update Compatibility On Startup" to get up-to-date information. + Display Compatibility Data:\nDisplays game compatibility information in table view. Enable "Update Compatibility On Startup" to get up-to-date information. + + + Update Compatibility On Startup:\nAutomatically update the compatibility database when shadPS4 starts. + Update Compatibility On Startup:\nAutomatically update the compatibility database when shadPS4 starts. + + + Update Compatibility Database:\nImmediately update the compatibility database. + Update Compatibility Database:\nImmediately update the compatibility database. + + + Never + Never + + + Idle + Idle + + + Always + Always + + + Touchpad Left + Touchpad Left + + + Touchpad Right + Touchpad Right + + + Touchpad Center + Touchpad Center + + + None + None + + + Graphics Device:\nOn multiple GPU systems, select the GPU the emulator will use from the drop down list,\nor select "Auto Select" to automatically determine it. + Graphics Device:\nOn multiple GPU systems, select the GPU the emulator will use from the drop down list,\nor select "Auto Select" to automatically determine it. + + + Width/Height:\nSets the size of the emulator window at launch, which can be resized during gameplay.\nThis is different from the in-game resolution. + Width/Height:\nSets the size of the emulator window at launch, which can be resized during gameplay.\nThis is different from the in-game resolution. + + + Vblank Divider:\nThe frame rate at which the emulator refreshes at is multiplied by this number. Changing this may have adverse effects, such as increasing the game speed, or breaking critical game functionality that does not expect this to change! + Vblank Divider:\nThe frame rate at which the emulator refreshes at is multiplied by this number. Changing this may have adverse effects, such as increasing the game speed, or breaking critical game functionality that does not expect this to change! + + + Enable Shaders Dumping:\nFor the sake of technical debugging, saves the games shaders to a folder as they render. + Enable Shaders Dumping:\nFor the sake of technical debugging, saves the games shaders to a folder as they render. + + + Enable Null GPU:\nFor the sake of technical debugging, disables game rendering as if there were no graphics card. + Enable Null GPU:\nFor the sake of technical debugging, disables game rendering as if there were no graphics card. + + + Enable HDR:\nEnables HDR in games that support it.\nYour monitor must have support for the BT2020 PQ color space and the RGB10A2 swapchain format. + Enable HDR:\nEnables HDR in games that support it.\nYour monitor must have support for the BT2020 PQ color space and the RGB10A2 swapchain format. + + + Game Folders:\nThe list of folders to check for installed games. + Game Folders:\nThe list of folders to check for installed games. + + + Add:\nAdd a folder to the list. + Add:\nAdd a folder to the list. + + + Remove:\nRemove a folder from the list. + Remove:\nRemove a folder from the list. + + + Enable Debug Dumping:\nSaves the import and export symbols and file header information of the currently running PS4 program to a directory. + Enable Debug Dumping:\nSaves the import and export symbols and file header information of the currently running PS4 program to a directory. + + + Enable Vulkan Validation Layers:\nEnables a system that validates the state of the Vulkan renderer and logs information about its internal state.\nThis will reduce performance and likely change the behavior of emulation. + Enable Vulkan Validation Layers:\nEnables a system that validates the state of the Vulkan renderer and logs information about its internal state.\nThis will reduce performance and likely change the behavior of emulation. + + + Enable Vulkan Synchronization Validation:\nEnables a system that validates the timing of Vulkan rendering tasks.\nThis will reduce performance and likely change the behavior of emulation. + Enable Vulkan Synchronization Validation:\nEnables a system that validates the timing of Vulkan rendering tasks.\nThis will reduce performance and likely change the behavior of emulation. + + + Enable RenderDoc Debugging:\nIf enabled, the emulator will provide compatibility with Renderdoc to allow capture and analysis of the currently rendered frame. + Enable RenderDoc Debugging:\nIf enabled, the emulator will provide compatibility with Renderdoc to allow capture and analysis of the currently rendered frame. + + + Collect Shaders:\nYou need this enabled to edit shaders with the debug menu (Ctrl + F10). + Collect Shaders:\nYou need this enabled to edit shaders with the debug menu (Ctrl + F10). + + + Crash Diagnostics:\nCreates a .yaml file with info about the Vulkan state at the time of crashing.\nUseful for debugging 'Device lost' errors. If you have this enabled, you should enable Host AND Guest Debug Markers.\nDoes not work on Intel GPUs.\nYou need Vulkan Validation Layers enabled and the Vulkan SDK for this to work. + Crash Diagnostics:\nCreates a .yaml file with info about the Vulkan state at the time of crashing.\nUseful for debugging 'Device lost' errors. If you have this enabled, you should enable Host AND Guest Debug Markers.\nDoes not work on Intel GPUs.\nYou need Vulkan Validation Layers enabled and the Vulkan SDK for this to work. + + + Copy GPU Buffers:\nGets around race conditions involving GPU submits.\nMay or may not help with PM4 type 0 crashes. + Copy GPU Buffers:\nGets around race conditions involving GPU submits.\nMay or may not help with PM4 type 0 crashes. + + + Host Debug Markers:\nInserts emulator-side information like markers for specific AMDGPU commands around Vulkan commands, as well as giving resources debug names.\nIf you have this enabled, you should enable Crash Diagnostics.\nUseful for programs like RenderDoc. + Host Debug Markers:\nInserts emulator-side information like markers for specific AMDGPU commands around Vulkan commands, as well as giving resources debug names.\nIf you have this enabled, you should enable Crash Diagnostics.\nUseful for programs like RenderDoc. + + + Guest Debug Markers:\nInserts any debug markers the game itself has added to the command buffer.\nIf you have this enabled, you should enable Crash Diagnostics.\nUseful for programs like RenderDoc. + Guest Debug Markers:\nInserts any debug markers the game itself has added to the command buffer.\nIf you have this enabled, you should enable Crash Diagnostics.\nUseful for programs like RenderDoc. + + + Save Data Path:\nThe folder where game save data will be saved. + Save Data Path:\nThe folder where game save data will be saved. + + + Browse:\nBrowse for a folder to set as the save data path. + Browse:\nBrowse for a folder to set as the save data path. + + + Borderless + Borderless + + + True + True + + + Release + Release + + + Nightly + Nightly + + + Set the volume of the background music. + Set the volume of the background music. + + + Enable Motion Controls + Enable Motion Controls + + + Save Data Path + Save Data Path + + + Browse + Browse + + + async + async + + + sync + sync + + + Auto Select + Auto Select + + + Directory to install games + Directory to install games + + + Directory to save data + Directory to save data + + + + TrophyViewer + + Trophy Viewer + Trophy Viewer + + + diff --git a/src/qt_gui/translations/ru_RU.ts b/src/qt_gui/translations/ru_RU.ts index 157fbd4cb..538b774fc 100644 --- a/src/qt_gui/translations/ru_RU.ts +++ b/src/qt_gui/translations/ru_RU.ts @@ -471,15 +471,15 @@ KBM Controls - KBM Controls + Управление KBM KBM Editor - KBM Editor + Редактор KBM Back - Back + Назад R1 / RB @@ -1238,7 +1238,7 @@ Package - Package + Пакет diff --git a/src/qt_gui/translations/tr_TR.ts b/src/qt_gui/translations/tr_TR.ts index b1aeaa9c3..34baf29bd 100644 --- a/src/qt_gui/translations/tr_TR.ts +++ b/src/qt_gui/translations/tr_TR.ts @@ -443,7 +443,7 @@ Left Deadzone - Left Deadzone + Sol Ölü Bölge Left Stick @@ -451,7 +451,7 @@ Config Selection - Config Selection + Yapılandırma Seçimi Common Config @@ -459,7 +459,7 @@ Use per-game configs - Use per-game configs + Oyuna özel yapılandırmaları kullan L1 / LB @@ -495,7 +495,7 @@ Options / Start - Options / Start + Seçenekler / Başlat R3 @@ -503,7 +503,7 @@ Face Buttons - Face Buttons + Eylem Düğmeleri Triangle / Y @@ -527,7 +527,7 @@ Right Deadzone - Right Deadzone + Sağ Ölü Bölge Right Stick @@ -655,7 +655,7 @@ Game has game-breaking glitches or unplayable performance - Game has game-breaking glitches or unplayable performance + Oyunu bozan hatalar ya da oynanamayan performans Game can be completed with playable performance and no major glitches @@ -725,7 +725,7 @@ Open Log Folder - Log Klasörünü Aç + Günlük Klasörünü Aç Copy info... @@ -809,7 +809,7 @@ This game has no update to delete! - This game has no update to delete! + Bu oyunun silinecek güncellemesi yok! Update @@ -817,7 +817,7 @@ This game has no DLC to delete! - This game has no DLC to delete! + Bu oyunun silinecek indirilebilir içeriği yok! DLC @@ -841,7 +841,7 @@ This game has no update folder to open! - This game has no update folder to open! + Bu oyunun açılacak güncelleme klasörü yok! Failed to convert icon. @@ -849,7 +849,7 @@ This game has no save data to delete! - This game has no save data to delete! + Bu oyunun silinecek kayıt verisi yok! Save Data @@ -1206,15 +1206,15 @@ Type - Type + Tür App Ver - App Ver + Uygulama Sürümü FW - FW + Sistem Yazılımı Region @@ -1349,7 +1349,7 @@ Back Button Behavior - Geri Dön Butonu Davranışı + Geri Dönme Butonu Davranışı Graphics @@ -1437,19 +1437,19 @@ Collect Shaders - Collect Shaders + Gölgelendiricileri Topla Copy GPU Buffers - Copy GPU Buffers + GPU Arabelleklerini Kopyala Host Debug Markers - Host Debug Markers + Ana Bilgisayar Hata Ayıklama İşaretleyicileri Guest Debug Markers - Guest Debug Markers + Konuk Hata Ayıklama İşaretleyicileri Update @@ -1569,7 +1569,7 @@ Trophy Key:\nKey used to decrypt trophies. Must be obtained from your jailbroken console.\nMust contain only hex characters. - Trophy Key:\nKey used to decrypt trophies. Must be obtained from your jailbroken console.\nMust contain only hex characters. + Kupa Anahtarı:\nKupaların şifresini çözmek için kullanılan anahtardır. Jailbreak yapılmış konsolunuzdan alınmalıdır.\nYalnızca onaltılık karakterler içermelidir. Log Type:\nSets whether to synchronize the output of the log window for performance. May have adverse effects on emulation. @@ -1585,11 +1585,11 @@ Background Image:\nControl the opacity of the game background image. - Background Image:\nControl the opacity of the game background image. + Arka Plan Resmi:\nOyunun arka plan resmi görünürlüğünü ayarlayın. Play Title Music:\nIf a game supports it, enable playing special music when selecting the game in the GUI. - Başlık Müziklerini Çal:\nEğer bir oyun bunu destekliyorsa, GUI'de oyunu seçtiğinizde özel müziklerin çalmasını etkinleştirir. + Oyun Müziklerini Çal:\nEğer oyun destekliyorsa, arayüzde oyunu seçtiğinizde özel müzik çalmasını etkinleştirir. Disable Trophy Pop-ups:\nDisable in-game trophy notifications. Trophy progress can still be tracked using the Trophy Viewer (right-click the game in the main window). @@ -1613,11 +1613,11 @@ Update Compatibility On Startup:\nAutomatically update the compatibility database when shadPS4 starts. - Update Compatibility On Startup:\nAutomatically update the compatibility database when shadPS4 starts. + Başlangıçta Uyumluluk Veritabanını Güncelle:\nshadPS4 başlatıldığında uyumluluk veritabanını otomatik olarak güncelleyin. Update Compatibility Database:\nImmediately update the compatibility database. - Update Compatibility Database:\nImmediately update the compatibility database. + Uyumluluk Veritabanını Güncelle:\nUyumluluk veri tabanını hemen güncelleyin. Never @@ -1721,7 +1721,7 @@ Save Data Path:\nThe folder where game save data will be saved. - Save Data Path:\nThe folder where game save data will be saved. + Kayıt Verileri Yolu:\nOyun kayıt verilerinin kaydedileceği klasördür. Browse:\nBrowse for a folder to set as the save data path. diff --git a/src/qt_gui/translations/zh_CN.ts b/src/qt_gui/translations/zh_CN.ts index 8eae7ae69..21daafb5e 100644 --- a/src/qt_gui/translations/zh_CN.ts +++ b/src/qt_gui/translations/zh_CN.ts @@ -495,7 +495,7 @@ Options / Start - 选项 / 开始 + Options / Start R3 @@ -1238,7 +1238,7 @@ Package - Package + Package @@ -1377,7 +1377,7 @@ Vblank Divider - Vblank Divider + Vblank Divider Advanced diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 3712380f5..2a5b9335e 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -242,14 +242,17 @@ void SetupCapabilities(const Info& info, const Profile& profile, EmitContext& ct ctx.AddCapability(spv::Capability::Image1D); ctx.AddCapability(spv::Capability::Sampled1D); ctx.AddCapability(spv::Capability::ImageQuery); + ctx.AddCapability(spv::Capability::Int8); + ctx.AddCapability(spv::Capability::Int16); + ctx.AddCapability(spv::Capability::Int64); + ctx.AddCapability(spv::Capability::UniformAndStorageBuffer8BitAccess); + ctx.AddCapability(spv::Capability::UniformAndStorageBuffer16BitAccess); if (info.uses_fp16) { ctx.AddCapability(spv::Capability::Float16); - ctx.AddCapability(spv::Capability::Int16); } if (info.uses_fp64) { ctx.AddCapability(spv::Capability::Float64); } - ctx.AddCapability(spv::Capability::Int64); if (info.has_storage_images) { ctx.AddCapability(spv::Capability::StorageImageExtendedFormats); ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat); diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp index ce65a5ccb..92cfcbb0f 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_atomic.cpp @@ -23,10 +23,13 @@ Id SharedAtomicU32(EmitContext& ctx, Id offset, Id value, Id BufferAtomicU32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value, Id (Sirit::Module::*atomic_func)(Id, Id, Id, Id, Id)) { - auto& buffer = ctx.buffers[handle]; - address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset); + const auto& buffer = ctx.buffers[handle]; + if (Sirit::ValidId(buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset); + } const Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(2u)); - const Id ptr = ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index); + const auto [id, pointer_type] = buffer[EmitContext::BufferAlias::U32]; + const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index); const auto [scope, semantics]{AtomicArgs(ctx)}; return (ctx.*atomic_func)(ctx.U32[1], ptr, scope, semantics, value); } @@ -165,17 +168,17 @@ Id EmitImageAtomicExchange32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id co } Id EmitDataAppend(EmitContext& ctx, u32 gds_addr, u32 binding) { - auto& buffer = ctx.buffers[binding]; - const Id ptr = ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, - ctx.ConstU32(gds_addr)); + const auto& buffer = ctx.buffers[binding]; + const auto [id, pointer_type] = buffer[EmitContext::BufferAlias::U32]; + const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, ctx.ConstU32(gds_addr)); const auto [scope, semantics]{AtomicArgs(ctx)}; return ctx.OpAtomicIIncrement(ctx.U32[1], ptr, scope, semantics); } Id EmitDataConsume(EmitContext& ctx, u32 gds_addr, u32 binding) { - auto& buffer = ctx.buffers[binding]; - const Id ptr = ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, - ctx.ConstU32(gds_addr)); + const auto& buffer = ctx.buffers[binding]; + const auto [id, pointer_type] = buffer[EmitContext::BufferAlias::U32]; + const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, ctx.ConstU32(gds_addr)); const auto [scope, semantics]{AtomicArgs(ctx)}; return ctx.OpAtomicIDecrement(ctx.U32[1], ptr, scope, semantics); } diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp index ae77ed413..cc7b7e097 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp @@ -160,21 +160,25 @@ void EmitGetGotoVariable(EmitContext&) { UNREACHABLE_MSG("Unreachable instruction"); } +using BufferAlias = EmitContext::BufferAlias; + Id EmitReadConst(EmitContext& ctx, IR::Inst* inst) { - u32 flatbuf_off_dw = inst->Flags(); - ASSERT(ctx.srt_flatbuf.binding >= 0); - ASSERT(flatbuf_off_dw > 0); - Id index = ctx.ConstU32(flatbuf_off_dw); - auto& buffer = ctx.srt_flatbuf; - const Id ptr{ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index)}; + const u32 flatbuf_off_dw = inst->Flags(); + const auto& srt_flatbuf = ctx.buffers.back(); + ASSERT(srt_flatbuf.binding >= 0 && flatbuf_off_dw > 0 && + srt_flatbuf.buffer_type == BufferType::ReadConstUbo); + const auto [id, pointer_type] = srt_flatbuf[BufferAlias::U32]; + const Id ptr{ + ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, ctx.ConstU32(flatbuf_off_dw))}; return ctx.OpLoad(ctx.U32[1], ptr); } Id EmitReadConstBuffer(EmitContext& ctx, u32 handle, Id index) { - auto& buffer = ctx.buffers[handle]; + const auto& buffer = ctx.buffers[handle]; index = ctx.OpIAdd(ctx.U32[1], index, buffer.offset_dwords); - const Id ptr{ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index)}; - return ctx.OpLoad(buffer.data_types->Get(1), ptr); + const auto [id, pointer_type] = buffer[BufferAlias::U32]; + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index)}; + return ctx.OpLoad(ctx.U32[1], ptr); } Id EmitReadStepRate(EmitContext& ctx, int rate_idx) { @@ -184,7 +188,7 @@ Id EmitReadStepRate(EmitContext& ctx, int rate_idx) { rate_idx == 0 ? ctx.u32_zero_value : ctx.u32_one_value)); } -Id EmitGetAttributeForGeometry(EmitContext& ctx, IR::Attribute attr, u32 comp, Id index) { +static Id EmitGetAttributeForGeometry(EmitContext& ctx, IR::Attribute attr, u32 comp, Id index) { if (IR::IsPosition(attr)) { ASSERT(attr == IR::Attribute::Position0); const auto position_arr_ptr = ctx.TypePointer(spv::StorageClass::Input, ctx.F32[4]); @@ -285,6 +289,8 @@ Id EmitGetAttributeU32(EmitContext& ctx, IR::Attribute attr, u32 comp) { return EmitReadStepRate(ctx, 0); case IR::Attribute::InstanceId1: return EmitReadStepRate(ctx, 1); + case IR::Attribute::WorkgroupIndex: + return ctx.workgroup_index_id; case IR::Attribute::WorkgroupId: return ctx.OpCompositeExtract(ctx.U32[1], ctx.OpLoad(ctx.U32[3], ctx.workgroup_id), comp); case IR::Attribute::LocalInvocationId: @@ -396,140 +402,158 @@ void EmitSetPatch(EmitContext& ctx, IR::Patch patch, Id value) { ctx.OpStore(pointer, value); } -template -static Id EmitLoadBufferU32xN(EmitContext& ctx, u32 handle, Id address) { - auto& buffer = ctx.buffers[handle]; - address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset); +template +static Id EmitLoadBufferB32xN(EmitContext& ctx, u32 handle, Id address) { + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } const Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(2u)); + const auto& data_types = alias == BufferAlias::U32 ? ctx.U32 : ctx.F32; + const auto [id, pointer_type] = spv_buffer[alias]; if constexpr (N == 1) { - const Id ptr{ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index)}; - return ctx.OpLoad(buffer.data_types->Get(1), ptr); + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index)}; + return ctx.OpLoad(data_types[1], ptr); } else { boost::container::static_vector ids; for (u32 i = 0; i < N; i++) { const Id index_i = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(i)); - const Id ptr{ - ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index_i)}; - ids.push_back(ctx.OpLoad(buffer.data_types->Get(1), ptr)); + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i)}; + ids.push_back(ctx.OpLoad(data_types[1], ptr)); } - return ctx.OpCompositeConstruct(buffer.data_types->Get(N), ids); + return ctx.OpCompositeConstruct(data_types[N], ids); } } Id EmitLoadBufferU8(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - const Id byte_index{ctx.OpBitwiseAnd(ctx.U32[1], address, ctx.ConstU32(3u))}; - const Id bit_offset{ctx.OpShiftLeftLogical(ctx.U32[1], byte_index, ctx.ConstU32(3u))}; - const Id dword{EmitLoadBufferU32xN<1>(ctx, handle, address)}; - return ctx.OpBitFieldUExtract(ctx.U32[1], dword, bit_offset, ctx.ConstU32(8u)); + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } + const auto [id, pointer_type] = spv_buffer[BufferAlias::U8]; + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, address)}; + return ctx.OpUConvert(ctx.U32[1], ctx.OpLoad(ctx.U8, ptr)); } Id EmitLoadBufferU16(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - const Id byte_index{ctx.OpBitwiseAnd(ctx.U32[1], address, ctx.ConstU32(2u))}; - const Id bit_offset{ctx.OpShiftLeftLogical(ctx.U32[1], byte_index, ctx.ConstU32(3u))}; - const Id dword{EmitLoadBufferU32xN<1>(ctx, handle, address)}; - return ctx.OpBitFieldUExtract(ctx.U32[1], dword, bit_offset, ctx.ConstU32(16u)); + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } + const auto [id, pointer_type] = spv_buffer[BufferAlias::U16]; + const Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(1u)); + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index)}; + return ctx.OpUConvert(ctx.U32[1], ctx.OpLoad(ctx.U16, ptr)); } Id EmitLoadBufferU32(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - return EmitLoadBufferU32xN<1>(ctx, handle, address); + return EmitLoadBufferB32xN<1, BufferAlias::U32>(ctx, handle, address); } Id EmitLoadBufferU32x2(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - return EmitLoadBufferU32xN<2>(ctx, handle, address); + return EmitLoadBufferB32xN<2, BufferAlias::U32>(ctx, handle, address); } Id EmitLoadBufferU32x3(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - return EmitLoadBufferU32xN<3>(ctx, handle, address); + return EmitLoadBufferB32xN<3, BufferAlias::U32>(ctx, handle, address); } Id EmitLoadBufferU32x4(EmitContext& ctx, IR::Inst*, u32 handle, Id address) { - return EmitLoadBufferU32xN<4>(ctx, handle, address); + return EmitLoadBufferB32xN<4, BufferAlias::U32>(ctx, handle, address); } Id EmitLoadBufferF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) { - return ctx.OpBitcast(ctx.F32[1], EmitLoadBufferU32(ctx, inst, handle, address)); + return EmitLoadBufferB32xN<1, BufferAlias::F32>(ctx, handle, address); } Id EmitLoadBufferF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) { - return ctx.OpBitcast(ctx.F32[2], EmitLoadBufferU32x2(ctx, inst, handle, address)); + return EmitLoadBufferB32xN<2, BufferAlias::F32>(ctx, handle, address); } Id EmitLoadBufferF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) { - return ctx.OpBitcast(ctx.F32[3], EmitLoadBufferU32x3(ctx, inst, handle, address)); + return EmitLoadBufferB32xN<3, BufferAlias::F32>(ctx, handle, address); } Id EmitLoadBufferF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) { - return ctx.OpBitcast(ctx.F32[4], EmitLoadBufferU32x4(ctx, inst, handle, address)); + return EmitLoadBufferB32xN<4, BufferAlias::F32>(ctx, handle, address); } Id EmitLoadBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) { UNREACHABLE_MSG("SPIR-V instruction"); } -template -static void EmitStoreBufferU32xN(EmitContext& ctx, u32 handle, Id address, Id value) { - auto& buffer = ctx.buffers[handle]; - address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset); +template +static void EmitStoreBufferB32xN(EmitContext& ctx, u32 handle, Id address, Id value) { + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } const Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(2u)); + const auto& data_types = alias == BufferAlias::U32 ? ctx.U32 : ctx.F32; + const auto [id, pointer_type] = spv_buffer[alias]; if constexpr (N == 1) { - const Id ptr{ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index)}; + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index)}; ctx.OpStore(ptr, value); } else { for (u32 i = 0; i < N; i++) { const Id index_i = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(i)); - const Id ptr = - ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index_i); - ctx.OpStore(ptr, ctx.OpCompositeExtract(buffer.data_types->Get(1), value, i)); + const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i); + ctx.OpStore(ptr, ctx.OpCompositeExtract(data_types[1], value, i)); } } } void EmitStoreBufferU8(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - const Id byte_index{ctx.OpBitwiseAnd(ctx.U32[1], address, ctx.ConstU32(3u))}; - const Id bit_offset{ctx.OpShiftLeftLogical(ctx.U32[1], byte_index, ctx.ConstU32(3u))}; - const Id dword{EmitLoadBufferU32xN<1>(ctx, handle, address)}; - const Id new_val{ctx.OpBitFieldInsert(ctx.U32[1], dword, value, bit_offset, ctx.ConstU32(8u))}; - EmitStoreBufferU32xN<1>(ctx, handle, address, new_val); + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } + const auto [id, pointer_type] = spv_buffer[BufferAlias::U8]; + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, address)}; + ctx.OpStore(ptr, ctx.OpUConvert(ctx.U8, value)); } void EmitStoreBufferU16(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - const Id byte_index{ctx.OpBitwiseAnd(ctx.U32[1], address, ctx.ConstU32(2u))}; - const Id bit_offset{ctx.OpShiftLeftLogical(ctx.U32[1], byte_index, ctx.ConstU32(3u))}; - const Id dword{EmitLoadBufferU32xN<1>(ctx, handle, address)}; - const Id new_val{ctx.OpBitFieldInsert(ctx.U32[1], dword, value, bit_offset, ctx.ConstU32(16u))}; - EmitStoreBufferU32xN<1>(ctx, handle, address, new_val); + const auto& spv_buffer = ctx.buffers[handle]; + if (Sirit::ValidId(spv_buffer.offset)) { + address = ctx.OpIAdd(ctx.U32[1], address, spv_buffer.offset); + } + const auto [id, pointer_type] = spv_buffer[BufferAlias::U16]; + const Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(1u)); + const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index)}; + ctx.OpStore(ptr, ctx.OpUConvert(ctx.U16, value)); } void EmitStoreBufferU32(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - EmitStoreBufferU32xN<1>(ctx, handle, address, value); + EmitStoreBufferB32xN<1, BufferAlias::U32>(ctx, handle, address, value); } void EmitStoreBufferU32x2(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - EmitStoreBufferU32xN<2>(ctx, handle, address, value); + EmitStoreBufferB32xN<2, BufferAlias::U32>(ctx, handle, address, value); } void EmitStoreBufferU32x3(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - EmitStoreBufferU32xN<3>(ctx, handle, address, value); + EmitStoreBufferB32xN<3, BufferAlias::U32>(ctx, handle, address, value); } void EmitStoreBufferU32x4(EmitContext& ctx, IR::Inst*, u32 handle, Id address, Id value) { - EmitStoreBufferU32xN<4>(ctx, handle, address, value); + EmitStoreBufferB32xN<4, BufferAlias::U32>(ctx, handle, address, value); } void EmitStoreBufferF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) { - EmitStoreBufferU32(ctx, inst, handle, address, ctx.OpBitcast(ctx.U32[1], value)); + EmitStoreBufferB32xN<1, BufferAlias::F32>(ctx, handle, address, value); } void EmitStoreBufferF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) { - EmitStoreBufferU32x2(ctx, inst, handle, address, ctx.OpBitcast(ctx.U32[2], value)); + EmitStoreBufferB32xN<2, BufferAlias::F32>(ctx, handle, address, value); } void EmitStoreBufferF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) { - EmitStoreBufferU32x3(ctx, inst, handle, address, ctx.OpBitcast(ctx.U32[3], value)); + EmitStoreBufferB32xN<3, BufferAlias::F32>(ctx, handle, address, value); } void EmitStoreBufferF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) { - EmitStoreBufferU32x4(ctx, inst, handle, address, ctx.OpBitcast(ctx.U32[4], value)); + EmitStoreBufferB32xN<4, BufferAlias::F32>(ctx, handle, address, value); } void EmitStoreBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp index 550b95f3d..8b1610d61 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp @@ -9,65 +9,35 @@ namespace Shader::Backend::SPIRV { Id EmitLoadSharedU32(EmitContext& ctx, Id offset) { const Id shift_id{ctx.ConstU32(2U)}; const Id index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)}; - if (ctx.info.has_emulated_shared_memory) { - const Id pointer = - ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, ctx.u32_zero_value, index); - return ctx.OpLoad(ctx.U32[1], pointer); - } else { - const Id pointer = ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, index); - return ctx.OpLoad(ctx.U32[1], pointer); - } + const Id pointer = ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, index); + return ctx.OpLoad(ctx.U32[1], pointer); } Id EmitLoadSharedU64(EmitContext& ctx, Id offset) { const Id shift_id{ctx.ConstU32(2U)}; const Id base_index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)}; const Id next_index{ctx.OpIAdd(ctx.U32[1], base_index, ctx.ConstU32(1U))}; - if (ctx.info.has_emulated_shared_memory) { - const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, - ctx.u32_zero_value, base_index)}; - const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, - ctx.u32_zero_value, next_index)}; - return ctx.OpCompositeConstruct(ctx.U32[2], ctx.OpLoad(ctx.U32[1], lhs_pointer), - ctx.OpLoad(ctx.U32[1], rhs_pointer)); - } else { - const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, base_index)}; - const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, next_index)}; - return ctx.OpCompositeConstruct(ctx.U32[2], ctx.OpLoad(ctx.U32[1], lhs_pointer), - ctx.OpLoad(ctx.U32[1], rhs_pointer)); - } + const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, base_index)}; + const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, next_index)}; + return ctx.OpCompositeConstruct(ctx.U32[2], ctx.OpLoad(ctx.U32[1], lhs_pointer), + ctx.OpLoad(ctx.U32[1], rhs_pointer)); } void EmitWriteSharedU32(EmitContext& ctx, Id offset, Id value) { const Id shift{ctx.ConstU32(2U)}; const Id word_offset{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift)}; - if (ctx.info.has_emulated_shared_memory) { - const Id pointer = ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, - ctx.u32_zero_value, word_offset); - ctx.OpStore(pointer, value); - } else { - const Id pointer = ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, word_offset); - ctx.OpStore(pointer, value); - } + const Id pointer = ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, word_offset); + ctx.OpStore(pointer, value); } void EmitWriteSharedU64(EmitContext& ctx, Id offset, Id value) { const Id shift{ctx.ConstU32(2U)}; const Id word_offset{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift)}; const Id next_offset{ctx.OpIAdd(ctx.U32[1], word_offset, ctx.ConstU32(1U))}; - if (ctx.info.has_emulated_shared_memory) { - const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, - ctx.u32_zero_value, word_offset)}; - const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, - ctx.u32_zero_value, next_offset)}; - ctx.OpStore(lhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 0U)); - ctx.OpStore(rhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 1U)); - } else { - const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, word_offset)}; - const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, next_offset)}; - ctx.OpStore(lhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 0U)); - ctx.OpStore(rhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 1U)); - } + const Id lhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, word_offset)}; + const Id rhs_pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, next_offset)}; + ctx.OpStore(lhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 0U)); + ctx.OpStore(rhs_pointer, ctx.OpCompositeExtract(ctx.U32[1], value, 1U)); } } // namespace Shader::Backend::SPIRV diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp index a0a3ed8ff..724550cd6 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_special.cpp @@ -11,6 +11,9 @@ void EmitPrologue(EmitContext& ctx) { if (ctx.stage == Stage::Fragment) { ctx.DefineInterpolatedAttribs(); } + if (ctx.info.loads.Get(IR::Attribute::WorkgroupIndex)) { + ctx.DefineWorkgroupIndex(); + } ctx.DefineBufferOffsets(); } diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index d676d205d..da20dc691 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -5,7 +5,6 @@ #include "common/div_ceil.h" #include "shader_recompiler/backend/spirv/spirv_emit_context.h" #include "shader_recompiler/frontend/fetch_shader.h" -#include "shader_recompiler/ir/passes/srt.h" #include "shader_recompiler/runtime_info.h" #include "video_core/amdgpu/types.h" @@ -107,6 +106,8 @@ Id EmitContext::Def(const IR::Value& value) { void EmitContext::DefineArithmeticTypes() { void_id = Name(TypeVoid(), "void_id"); U1[1] = Name(TypeBool(), "bool_id"); + U8 = Name(TypeUInt(8), "u8_id"); + U16 = Name(TypeUInt(16), "u16_id"); if (info.uses_fp16) { F16[1] = Name(TypeFloat(16), "f16_id"); U16 = Name(TypeUInt(16), "u16_id"); @@ -193,6 +194,9 @@ EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat f void EmitContext::DefineBufferOffsets() { for (BufferDefinition& buffer : buffers) { + if (buffer.buffer_type != BufferType::Guest) { + continue; + } const u32 binding = buffer.binding; const u32 half = PushData::BufOffsetIndex + (binding >> 4); const u32 comp = (binding & 0xf) >> 2; @@ -211,8 +215,7 @@ void EmitContext::DefineInterpolatedAttribs() { if (!profile.needs_manual_interpolation) { return; } - // Iterate all input attributes, load them and manually interpolate with barycentric - // coordinates. + // Iterate all input attributes, load them and manually interpolate. for (s32 i = 0; i < runtime_info.fs_info.num_inputs; i++) { const auto& input = runtime_info.fs_info.inputs[i]; const u32 semantic = input.param_index; @@ -237,6 +240,20 @@ void EmitContext::DefineInterpolatedAttribs() { } } +void EmitContext::DefineWorkgroupIndex() { + const Id workgroup_id_val{OpLoad(U32[3], workgroup_id)}; + const Id workgroup_x{OpCompositeExtract(U32[1], workgroup_id_val, 0)}; + const Id workgroup_y{OpCompositeExtract(U32[1], workgroup_id_val, 1)}; + const Id workgroup_z{OpCompositeExtract(U32[1], workgroup_id_val, 2)}; + const Id num_workgroups{OpLoad(U32[3], num_workgroups_id)}; + const Id num_workgroups_x{OpCompositeExtract(U32[1], num_workgroups, 0)}; + const Id num_workgroups_y{OpCompositeExtract(U32[1], num_workgroups, 1)}; + workgroup_index_id = + OpIAdd(U32[1], OpIAdd(U32[1], workgroup_x, OpIMul(U32[1], workgroup_y, num_workgroups_x)), + OpIMul(U32[1], workgroup_z, OpIMul(U32[1], num_workgroups_x, num_workgroups_y))); + Name(workgroup_index_id, "workgroup_index"); +} + Id MakeDefaultValue(EmitContext& ctx, u32 default_value) { switch (default_value) { case 0: @@ -305,9 +322,16 @@ void EmitContext::DefineInputs() { break; } case LogicalStage::Fragment: - frag_coord = DefineVariable(F32[4], spv::BuiltIn::FragCoord, spv::StorageClass::Input); - frag_depth = DefineVariable(F32[1], spv::BuiltIn::FragDepth, spv::StorageClass::Output); - front_facing = DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input); + if (info.loads.GetAny(IR::Attribute::FragCoord)) { + frag_coord = DefineVariable(F32[4], spv::BuiltIn::FragCoord, spv::StorageClass::Input); + } + if (info.stores.Get(IR::Attribute::Depth)) { + frag_depth = DefineVariable(F32[1], spv::BuiltIn::FragDepth, spv::StorageClass::Output); + } + if (info.loads.Get(IR::Attribute::IsFrontFace)) { + front_facing = + DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input); + } if (profile.needs_manual_interpolation) { gl_bary_coord_id = DefineVariable(F32[3], spv::BuiltIn::BaryCoordKHR, spv::StorageClass::Input); @@ -342,9 +366,19 @@ void EmitContext::DefineInputs() { } break; case LogicalStage::Compute: - workgroup_id = DefineVariable(U32[3], spv::BuiltIn::WorkgroupId, spv::StorageClass::Input); - local_invocation_id = - DefineVariable(U32[3], spv::BuiltIn::LocalInvocationId, spv::StorageClass::Input); + if (info.loads.GetAny(IR::Attribute::WorkgroupIndex) || + info.loads.GetAny(IR::Attribute::WorkgroupId)) { + workgroup_id = + DefineVariable(U32[3], spv::BuiltIn::WorkgroupId, spv::StorageClass::Input); + } + if (info.loads.GetAny(IR::Attribute::WorkgroupIndex)) { + num_workgroups_id = + DefineVariable(U32[3], spv::BuiltIn::NumWorkgroups, spv::StorageClass::Input); + } + if (info.loads.GetAny(IR::Attribute::LocalInvocationId)) { + local_invocation_id = + DefineVariable(U32[3], spv::BuiltIn::LocalInvocationId, spv::StorageClass::Input); + } break; case LogicalStage::Geometry: { primitive_id = DefineVariable(U32[1], spv::BuiltIn::PrimitiveId, spv::StorageClass::Input); @@ -588,78 +622,74 @@ void EmitContext::DefinePushDataBlock() { interfaces.push_back(push_data_block); } -void EmitContext::DefineBuffers() { - boost::container::small_vector type_ids; - const auto define_struct = [&](Id record_array_type, bool is_instance_data, - std::optional explicit_name = {}) { - const Id struct_type{TypeStruct(record_array_type)}; - if (std::ranges::find(type_ids, record_array_type.value, &Id::value) != type_ids.end()) { - return struct_type; - } - Decorate(record_array_type, spv::Decoration::ArrayStride, 4); - auto name = is_instance_data ? fmt::format("{}_instance_data_f32", stage) - : fmt::format("{}_cbuf_block_f32", stage); - name = explicit_name.value_or(name); - Name(struct_type, name); +EmitContext::BufferSpv EmitContext::DefineBuffer(bool is_storage, bool is_written, u32 elem_shift, + BufferType buffer_type, Id data_type) { + // Define array type. + const Id max_num_items = ConstU32(u32(profile.max_ubo_size) >> elem_shift); + const Id record_array_type{is_storage ? TypeRuntimeArray(data_type) + : TypeArray(data_type, max_num_items)}; + // Define block struct type. Don't perform decorations twice on the same Id. + const Id struct_type{TypeStruct(record_array_type)}; + if (std::ranges::find(buf_type_ids, record_array_type.value, &Id::value) == + buf_type_ids.end()) { + Decorate(record_array_type, spv::Decoration::ArrayStride, 1 << elem_shift); Decorate(struct_type, spv::Decoration::Block); MemberName(struct_type, 0, "data"); MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); - type_ids.push_back(record_array_type); - return struct_type; - }; - - if (info.has_readconst) { - const Id data_type = U32[1]; - const auto storage_class = spv::StorageClass::Uniform; - const Id pointer_type = TypePointer(storage_class, data_type); - const Id record_array_type{ - TypeArray(U32[1], ConstU32(static_cast(info.flattened_ud_buf.size())))}; - - const Id struct_type{define_struct(record_array_type, false, "srt_flatbuf_ty")}; - - const Id struct_pointer_type{TypePointer(storage_class, struct_type)}; - const Id id{AddGlobalVariable(struct_pointer_type, storage_class)}; - Decorate(id, spv::Decoration::Binding, binding.unified++); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - Name(id, "srt_flatbuf_ubo"); - - srt_flatbuf = { - .id = id, - .binding = binding.buffer++, - .pointer_type = pointer_type, - }; - interfaces.push_back(id); + buf_type_ids.push_back(record_array_type); } + // Define buffer binding interface. + const auto storage_class = + is_storage ? spv::StorageClass::StorageBuffer : spv::StorageClass::Uniform; + const Id struct_pointer_type{TypePointer(storage_class, struct_type)}; + const Id pointer_type = TypePointer(storage_class, data_type); + const Id id{AddGlobalVariable(struct_pointer_type, storage_class)}; + Decorate(id, spv::Decoration::Binding, binding.unified); + Decorate(id, spv::Decoration::DescriptorSet, 0U); + if (is_storage && !is_written) { + Decorate(id, spv::Decoration::NonWritable); + } + switch (buffer_type) { + case Shader::BufferType::GdsBuffer: + Name(id, "gds_buffer"); + break; + case Shader::BufferType::ReadConstUbo: + Name(id, "srt_flatbuf_ubo"); + break; + case Shader::BufferType::SharedMemory: + Name(id, "ssbo_shmem"); + break; + default: + Name(id, fmt::format("{}_{}", is_storage ? "ssbo" : "ubo", binding.buffer)); + } + interfaces.push_back(id); + return {id, pointer_type}; +}; +void EmitContext::DefineBuffers() { for (const auto& desc : info.buffers) { - const auto sharp = desc.GetSharp(info); - const bool is_storage = desc.IsStorage(sharp, profile); - const u32 array_size = profile.max_ubo_size >> 2; - const auto* data_types = True(desc.used_types & IR::Type::F32) ? &F32 : &U32; - const Id data_type = (*data_types)[1]; - const Id record_array_type{is_storage ? TypeRuntimeArray(data_type) - : TypeArray(data_type, ConstU32(array_size))}; - const Id struct_type{define_struct(record_array_type, desc.is_instance_data)}; + const auto buf_sharp = desc.GetSharp(info); + const bool is_storage = desc.IsStorage(buf_sharp, profile); - const auto storage_class = - is_storage ? spv::StorageClass::StorageBuffer : spv::StorageClass::Uniform; - const Id struct_pointer_type{TypePointer(storage_class, struct_type)}; - const Id pointer_type = TypePointer(storage_class, data_type); - const Id id{AddGlobalVariable(struct_pointer_type, storage_class)}; - Decorate(id, spv::Decoration::Binding, binding.unified++); - Decorate(id, spv::Decoration::DescriptorSet, 0U); - if (is_storage && !desc.is_written) { - Decorate(id, spv::Decoration::NonWritable); + // Define aliases depending on the shader usage. + auto& spv_buffer = buffers.emplace_back(binding.buffer++, desc.buffer_type); + if (True(desc.used_types & IR::Type::U32)) { + spv_buffer[BufferAlias::U32] = + DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, U32[1]); } - Name(id, fmt::format("{}_{}", is_storage ? "ssbo" : "cbuf", desc.sharp_idx)); - - buffers.push_back({ - .id = id, - .binding = binding.buffer++, - .data_types = data_types, - .pointer_type = pointer_type, - }); - interfaces.push_back(id); + if (True(desc.used_types & IR::Type::F32)) { + spv_buffer[BufferAlias::F32] = + DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, F32[1]); + } + if (True(desc.used_types & IR::Type::U16)) { + spv_buffer[BufferAlias::U16] = + DefineBuffer(is_storage, desc.is_written, 1, desc.buffer_type, U16); + } + if (True(desc.used_types & IR::Type::U8)) { + spv_buffer[BufferAlias::U8] = + DefineBuffer(is_storage, desc.is_written, 0, desc.buffer_type, U8); + } + ++binding.unified; } } @@ -809,51 +839,18 @@ void EmitContext::DefineImagesAndSamplers() { } void EmitContext::DefineSharedMemory() { - static constexpr size_t DefaultSharedMemSize = 2_KB; if (!info.uses_shared) { return; } ASSERT(info.stage == Stage::Compute); - - const u32 max_shared_memory_size = profile.max_shared_memory_size; - u32 shared_memory_size = runtime_info.cs_info.shared_memory_size; - if (shared_memory_size == 0) { - shared_memory_size = DefaultSharedMemSize; - } - + const u32 shared_memory_size = runtime_info.cs_info.shared_memory_size; const u32 num_elements{Common::DivCeil(shared_memory_size, 4U)}; const Id type{TypeArray(U32[1], ConstU32(num_elements))}; - - if (shared_memory_size <= max_shared_memory_size) { - shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type); - shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]); - shared_memory_u32 = AddGlobalVariable(shared_memory_u32_type, spv::StorageClass::Workgroup); - Name(shared_memory_u32, "shared_mem"); - interfaces.push_back(shared_memory_u32); - } else { - shared_memory_u32_type = TypePointer(spv::StorageClass::StorageBuffer, type); - shared_u32 = TypePointer(spv::StorageClass::StorageBuffer, U32[1]); - - Decorate(type, spv::Decoration::ArrayStride, 4); - - const Id struct_type{TypeStruct(type)}; - Name(struct_type, "shared_memory_buf"); - Decorate(struct_type, spv::Decoration::Block); - MemberName(struct_type, 0, "data"); - MemberDecorate(struct_type, 0, spv::Decoration::Offset, 0U); - - const Id struct_pointer_type{TypePointer(spv::StorageClass::StorageBuffer, struct_type)}; - const Id ssbo_id{AddGlobalVariable(struct_pointer_type, spv::StorageClass::StorageBuffer)}; - Decorate(ssbo_id, spv::Decoration::Binding, binding.unified++); - Decorate(ssbo_id, spv::Decoration::DescriptorSet, 0U); - Name(ssbo_id, "shared_mem_ssbo"); - - shared_memory_u32 = ssbo_id; - - info.has_emulated_shared_memory = true; - info.shared_memory_size = shared_memory_size; - interfaces.push_back(ssbo_id); - } + shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type); + shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]); + shared_memory_u32 = AddGlobalVariable(shared_memory_u32_type, spv::StorageClass::Workgroup); + Name(shared_memory_u32, "shared_mem"); + interfaces.push_back(shared_memory_u32); } Id EmitContext::DefineFloat32ToUfloatM5(u32 mantissa_bits, const std::string_view name) { diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h index 23fca4212..0fe6e336c 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -8,7 +8,7 @@ #include "shader_recompiler/backend/bindings.h" #include "shader_recompiler/info.h" -#include "shader_recompiler/ir/program.h" +#include "shader_recompiler/ir/value.h" #include "shader_recompiler/profile.h" namespace Shader::Backend::SPIRV { @@ -45,6 +45,7 @@ public: void DefineBufferOffsets(); void DefineInterpolatedAttribs(); + void DefineWorkgroupIndex(); [[nodiscard]] Id DefineInput(Id type, std::optional location = std::nullopt, std::optional builtin = std::nullopt) { @@ -200,8 +201,10 @@ public: std::array patches{}; Id workgroup_id{}; + Id num_workgroups_id{}; + Id workgroup_index_id{}; Id local_invocation_id{}; - Id invocation_id{}; // for instanced geoshaders or output vertices within TCS patch + Id invocation_id{}; Id subgroup_local_invocation_id{}; Id image_u32{}; @@ -227,18 +230,38 @@ public: bool is_storage = false; }; - struct BufferDefinition { + enum class BufferAlias : u32 { + U8, + U16, + U32, + F32, + NumAlias, + }; + + struct BufferSpv { Id id; - Id offset; - Id offset_dwords; - u32 binding; - const VectorIds* data_types; Id pointer_type; }; + struct BufferDefinition { + u32 binding; + BufferType buffer_type; + Id offset; + Id offset_dwords; + std::array aliases; + + const BufferSpv& operator[](BufferAlias alias) const { + return aliases[u32(alias)]; + } + + BufferSpv& operator[](BufferAlias alias) { + return aliases[u32(alias)]; + } + }; + Bindings& binding; + boost::container::small_vector buf_type_ids; boost::container::small_vector buffers; - BufferDefinition srt_flatbuf; boost::container::small_vector images; boost::container::small_vector samplers; @@ -279,6 +302,9 @@ private: SpirvAttribute GetAttributeInfo(AmdGpu::NumberFormat fmt, Id id, u32 num_components, bool output); + BufferSpv DefineBuffer(bool is_storage, bool is_written, u32 elem_shift, BufferType buffer_type, + Id data_type); + Id DefineFloat32ToUfloatM5(u32 mantissa_bits, std::string_view name); Id DefineUfloatM5ToFloat32(u32 mantissa_bits, std::string_view name); }; diff --git a/src/shader_recompiler/frontend/translate/data_share.cpp b/src/shader_recompiler/frontend/translate/data_share.cpp index 62c0423dd..460f8913c 100644 --- a/src/shader_recompiler/frontend/translate/data_share.cpp +++ b/src/shader_recompiler/frontend/translate/data_share.cpp @@ -176,6 +176,13 @@ void Translator::DS_WRITE(int bit_size, bool is_signed, bool is_pair, bool strid const IR::U32 addr{ir.GetVectorReg(IR::VectorReg(inst.src[0].code))}; const IR::VectorReg data0{inst.src[1].code}; const IR::VectorReg data1{inst.src[2].code}; + const u32 offset = (inst.control.ds.offset1 << 8u) + inst.control.ds.offset0; + if (info.stage == Stage::Fragment) { + ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0, + "Unexpected shared memory offset alignment: {}", offset); + ir.SetVectorReg(GetScratchVgpr(offset), ir.GetVectorReg(data0)); + return; + } if (is_pair) { const u32 adj = (bit_size == 32 ? 4 : 8) * (stride64 ? 64 : 1); const IR::U32 addr0 = ir.IAdd(addr, ir.Imm32(u32(inst.control.ds.offset0 * adj))); @@ -195,14 +202,12 @@ void Translator::DS_WRITE(int bit_size, bool is_signed, bool is_pair, bool strid addr1); } } else if (bit_size == 64) { - const IR::U32 addr0 = ir.IAdd( - addr, ir.Imm32((u32(inst.control.ds.offset1) << 8u) + u32(inst.control.ds.offset0))); + const IR::U32 addr0 = ir.IAdd(addr, ir.Imm32(offset)); const IR::Value data = ir.CompositeConstruct(ir.GetVectorReg(data0), ir.GetVectorReg(data0 + 1)); ir.WriteShared(bit_size, data, addr0); } else { - const IR::U32 addr0 = ir.IAdd( - addr, ir.Imm32((u32(inst.control.ds.offset1) << 8u) + u32(inst.control.ds.offset0))); + const IR::U32 addr0 = ir.IAdd(addr, ir.Imm32(offset)); ir.WriteShared(bit_size, ir.GetVectorReg(data0), addr0); } } @@ -223,6 +228,13 @@ void Translator::DS_READ(int bit_size, bool is_signed, bool is_pair, bool stride const GcnInst& inst) { const IR::U32 addr{ir.GetVectorReg(IR::VectorReg(inst.src[0].code))}; IR::VectorReg dst_reg{inst.dst[0].code}; + const u32 offset = (inst.control.ds.offset1 << 8u) + inst.control.ds.offset0; + if (info.stage == Stage::Fragment) { + ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0, + "Unexpected shared memory offset alignment: {}", offset); + ir.SetVectorReg(dst_reg, ir.GetVectorReg(GetScratchVgpr(offset))); + return; + } if (is_pair) { // Pair loads are either 32 or 64-bit const u32 adj = (bit_size == 32 ? 4 : 8) * (stride64 ? 64 : 1); @@ -243,14 +255,12 @@ void Translator::DS_READ(int bit_size, bool is_signed, bool is_pair, bool stride ir.SetVectorReg(dst_reg++, IR::U32{ir.CompositeExtract(data1, 1)}); } } else if (bit_size == 64) { - const IR::U32 addr0 = ir.IAdd( - addr, ir.Imm32((u32(inst.control.ds.offset1) << 8u) + u32(inst.control.ds.offset0))); + const IR::U32 addr0 = ir.IAdd(addr, ir.Imm32(offset)); const IR::Value data = ir.LoadShared(bit_size, is_signed, addr0); ir.SetVectorReg(dst_reg, IR::U32{ir.CompositeExtract(data, 0)}); ir.SetVectorReg(dst_reg + 1, IR::U32{ir.CompositeExtract(data, 1)}); } else { - const IR::U32 addr0 = ir.IAdd( - addr, ir.Imm32((u32(inst.control.ds.offset1) << 8u) + u32(inst.control.ds.offset0))); + const IR::U32 addr0 = ir.IAdd(addr, ir.Imm32(offset)); const IR::U32 data = IR::U32{ir.LoadShared(bit_size, is_signed, addr0)}; ir.SetVectorReg(dst_reg, data); } diff --git a/src/shader_recompiler/frontend/translate/export.cpp b/src/shader_recompiler/frontend/translate/export.cpp index ece35093a..0abef2e81 100644 --- a/src/shader_recompiler/frontend/translate/export.cpp +++ b/src/shader_recompiler/frontend/translate/export.cpp @@ -7,7 +7,7 @@ namespace Shader::Gcn { -u32 SwizzleMrtComponent(const FragmentRuntimeInfo::PsColorBuffer& color_buffer, u32 comp) { +u32 SwizzleMrtComponent(const PsColorBuffer& color_buffer, u32 comp) { const auto [r, g, b, a] = color_buffer.swizzle; const std::array swizzle_array = {r, g, b, a}; const auto swizzled_comp_type = static_cast(swizzle_array[comp]); @@ -16,7 +16,7 @@ u32 SwizzleMrtComponent(const FragmentRuntimeInfo::PsColorBuffer& color_buffer, } void Translator::ExportMrtValue(IR::Attribute attribute, u32 comp, const IR::F32& value, - const FragmentRuntimeInfo::PsColorBuffer& color_buffer) { + const PsColorBuffer& color_buffer) { auto converted = ApplyWriteNumberConversion(ir, value, color_buffer.num_conversion); if (color_buffer.needs_unorm_fixup) { // FIXME: Fix-up for GPUs where float-to-unorm rounding is off from expected. diff --git a/src/shader_recompiler/frontend/translate/translate.cpp b/src/shader_recompiler/frontend/translate/translate.cpp index 7f5504663..7f1bcb33e 100644 --- a/src/shader_recompiler/frontend/translate/translate.cpp +++ b/src/shader_recompiler/frontend/translate/translate.cpp @@ -4,7 +4,6 @@ #include "common/config.h" #include "common/io_file.h" #include "common/path_util.h" -#include "shader_recompiler/exception.h" #include "shader_recompiler/frontend/fetch_shader.h" #include "shader_recompiler/frontend/translate/translate.h" #include "shader_recompiler/info.h" @@ -21,9 +20,14 @@ namespace Shader::Gcn { +static u32 next_vgpr_num; +static std::unordered_map vgpr_map; + Translator::Translator(IR::Block* block_, Info& info_, const RuntimeInfo& runtime_info_, const Profile& profile_) - : ir{*block_, block_->begin()}, info{info_}, runtime_info{runtime_info_}, profile{profile_} {} + : ir{*block_, block_->begin()}, info{info_}, runtime_info{runtime_info_}, profile{profile_} { + next_vgpr_num = vgpr_map.empty() ? runtime_info.num_allocated_vgprs : next_vgpr_num; +} void Translator::EmitPrologue() { ir.Prologue(); @@ -179,8 +183,21 @@ void Translator::EmitPrologue() { default: UNREACHABLE_MSG("Unknown shader stage"); } + + // Clear any scratch vgpr mappings for next shader. + vgpr_map.clear(); } +IR::VectorReg Translator::GetScratchVgpr(u32 offset) { + const auto [it, is_new] = vgpr_map.try_emplace(offset); + if (is_new) { + ASSERT_MSG(next_vgpr_num < 256, "Out of VGPRs"); + const auto new_vgpr = static_cast(next_vgpr_num++); + it->second = new_vgpr; + } + return it->second; +}; + template T Translator::GetSrc(const InstOperand& operand) { constexpr bool is_float = std::is_same_v; @@ -490,7 +507,6 @@ void Translator::EmitFetch(const GcnInst& inst) { info.buffers.push_back({ .sharp_idx = info.srt_info.ReserveSharp(attrib.sgpr_base, attrib.dword_offset, 4), .used_types = IR::Type::F32, - .is_instance_data = true, .instance_attrib = attrib.semantic, }); } diff --git a/src/shader_recompiler/frontend/translate/translate.h b/src/shader_recompiler/frontend/translate/translate.h index 287885854..563881a8e 100644 --- a/src/shader_recompiler/frontend/translate/translate.h +++ b/src/shader_recompiler/frontend/translate/translate.h @@ -309,7 +309,7 @@ private: const IR::F32& x_res, const IR::F32& y_res, const IR::F32& z_res); void ExportMrtValue(IR::Attribute attribute, u32 comp, const IR::F32& value, - const FragmentRuntimeInfo::PsColorBuffer& color_buffer); + const PsColorBuffer& color_buffer); void ExportMrtCompressed(IR::Attribute attribute, u32 idx, const IR::U32& value); void ExportMrtUncompressed(IR::Attribute attribute, u32 comp, const IR::F32& value); void ExportCompressed(IR::Attribute attribute, u32 idx, const IR::U32& value); @@ -317,6 +317,8 @@ private: void LogMissingOpcode(const GcnInst& inst); + IR::VectorReg GetScratchVgpr(u32 offset); + private: IR::IREmitter ir; Info& info; diff --git a/src/shader_recompiler/info.h b/src/shader_recompiler/info.h index 57d428a49..13f310cf8 100644 --- a/src/shader_recompiler/info.h +++ b/src/shader_recompiler/info.h @@ -2,7 +2,6 @@ // SPDX-License-Identifier: GPL-2.0-or-later #pragma once -#include #include #include #include @@ -19,7 +18,6 @@ #include "shader_recompiler/params.h" #include "shader_recompiler/profile.h" #include "shader_recompiler/runtime_info.h" -#include "video_core/amdgpu/liverpool.h" #include "video_core/amdgpu/resource.h" namespace Shader { @@ -37,21 +35,30 @@ enum class TextureType : u32 { }; constexpr u32 NUM_TEXTURE_TYPES = 7; +enum class BufferType : u32 { + Guest, + ReadConstUbo, + GdsBuffer, + SharedMemory, +}; + struct Info; struct BufferResource { u32 sharp_idx; IR::Type used_types; AmdGpu::Buffer inline_cbuf; - bool is_gds_buffer{}; - bool is_instance_data{}; + BufferType buffer_type; u8 instance_attrib{}; bool is_written{}; bool is_formatted{}; - [[nodiscard]] bool IsStorage(const AmdGpu::Buffer& buffer, - const Profile& profile) const noexcept { - return buffer.GetSize() > profile.max_ubo_size || is_written || is_gds_buffer; + bool IsSpecial() const noexcept { + return buffer_type != BufferType::Guest; + } + + bool IsStorage(const AmdGpu::Buffer& buffer, const Profile& profile) const noexcept { + return buffer.GetSize() > profile.max_ubo_size || is_written; } [[nodiscard]] constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept; @@ -193,10 +200,8 @@ struct Info { bool uses_unpack_10_11_11{}; bool stores_tess_level_outer{}; bool stores_tess_level_inner{}; - bool translation_failed{}; // indicates that shader has unsupported instructions - bool has_emulated_shared_memory{}; + bool translation_failed{}; bool has_readconst{}; - u32 shared_memory_size{}; u8 mrt_mask{0u}; bool has_fetch_shader{false}; u32 fetch_shader_sgpr_base{0u}; @@ -233,10 +238,8 @@ struct Info { } void AddBindings(Backend::Bindings& bnd) const { - const auto total_buffers = - buffers.size() + (has_readconst ? 1 : 0) + (has_emulated_shared_memory ? 1 : 0); - bnd.buffer += total_buffers; - bnd.unified += total_buffers + images.size() + samplers.size(); + bnd.buffer += buffers.size(); + bnd.unified += buffers.size() + images.size() + samplers.size(); bnd.user_data += ud_mask.NumRegs(); } @@ -283,14 +286,3 @@ constexpr AmdGpu::Image FMaskResource::GetSharp(const Info& info) const noexcept } } // namespace Shader - -template <> -struct fmt::formatter { - constexpr auto parse(format_parse_context& ctx) { - return ctx.begin(); - } - auto format(const Shader::Stage stage, format_context& ctx) const { - constexpr static std::array names = {"fs", "vs", "gs", "es", "hs", "ls", "cs"}; - return fmt::format_to(ctx.out(), "{}", names[static_cast(stage)]); - } -}; diff --git a/src/shader_recompiler/ir/attribute.h b/src/shader_recompiler/ir/attribute.h index bcb2b44a9..5117f5650 100644 --- a/src/shader_recompiler/ir/attribute.h +++ b/src/shader_recompiler/ir/attribute.h @@ -69,16 +69,17 @@ enum class Attribute : u64 { SampleIndex = 72, GlobalInvocationId = 73, WorkgroupId = 74, - LocalInvocationId = 75, - LocalInvocationIndex = 76, - FragCoord = 77, - InstanceId0 = 78, // step rate 0 - InstanceId1 = 79, // step rate 1 - InvocationId = 80, // TCS id in output patch and instanced geometry shader id - PatchVertices = 81, - TessellationEvaluationPointU = 82, - TessellationEvaluationPointV = 83, - PackedHullInvocationInfo = 84, // contains patch id within the VGT and invocation ID + WorkgroupIndex = 75, + LocalInvocationId = 76, + LocalInvocationIndex = 77, + FragCoord = 78, + InstanceId0 = 79, // step rate 0 + InstanceId1 = 80, // step rate 1 + InvocationId = 81, // TCS id in output patch and instanced geometry shader id + PatchVertices = 82, + TessellationEvaluationPointU = 83, + TessellationEvaluationPointV = 84, + PackedHullInvocationInfo = 85, // contains patch id within the VGT and invocation ID Max, }; diff --git a/src/shader_recompiler/ir/passes/ir_passes.h b/src/shader_recompiler/ir/passes/ir_passes.h index 3c98579a0..69628dbfd 100644 --- a/src/shader_recompiler/ir/passes/ir_passes.h +++ b/src/shader_recompiler/ir/passes/ir_passes.h @@ -20,12 +20,14 @@ void FlattenExtendedUserdataPass(IR::Program& program); void ResourceTrackingPass(IR::Program& program); void CollectShaderInfoPass(IR::Program& program); void LowerBufferFormatToRaw(IR::Program& program); -void LowerSharedMemToRegisters(IR::Program& program, const RuntimeInfo& runtime_info); void RingAccessElimination(const IR::Program& program, const RuntimeInfo& runtime_info, Stage stage); void TessellationPreprocess(IR::Program& program, RuntimeInfo& runtime_info); void HullShaderTransform(IR::Program& program, RuntimeInfo& runtime_info); void DomainShaderTransform(IR::Program& program, RuntimeInfo& runtime_info); -void SharedMemoryBarrierPass(IR::Program& program, const Profile& profile); +void SharedMemoryBarrierPass(IR::Program& program, const RuntimeInfo& runtime_info, + const Profile& profile); +void SharedMemoryToStoragePass(IR::Program& program, const RuntimeInfo& runtime_info, + const Profile& profile); } // namespace Shader::Optimization diff --git a/src/shader_recompiler/ir/passes/lower_shared_mem_to_registers.cpp b/src/shader_recompiler/ir/passes/lower_shared_mem_to_registers.cpp deleted file mode 100644 index 23963a991..000000000 --- a/src/shader_recompiler/ir/passes/lower_shared_mem_to_registers.cpp +++ /dev/null @@ -1,81 +0,0 @@ -// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project -// SPDX-License-Identifier: GPL-2.0-or-later - -#include - -#include "shader_recompiler/ir/ir_emitter.h" -#include "shader_recompiler/ir/program.h" - -namespace Shader::Optimization { - -static bool IsSharedMemoryInst(const IR::Inst& inst) { - const auto opcode = inst.GetOpcode(); - return opcode == IR::Opcode::LoadSharedU32 || opcode == IR::Opcode::LoadSharedU64 || - opcode == IR::Opcode::WriteSharedU32 || opcode == IR::Opcode::WriteSharedU64; -} - -static u32 GetSharedMemImmOffset(const IR::Inst& inst) { - const auto* address = inst.Arg(0).InstRecursive(); - ASSERT(address->GetOpcode() == IR::Opcode::IAdd32); - const auto ir_offset = address->Arg(1); - ASSERT_MSG(ir_offset.IsImmediate()); - const auto offset = ir_offset.U32(); - // Typical usage is the compiler spilling registers into shared memory, with 256 bytes between - // each register to account for 4 bytes per register times 64 threads per group. Ensure that - // this assumption holds, as if it does not this approach may need to be revised. - ASSERT_MSG(offset % 256 == 0, "Unexpected shared memory offset alignment: {}", offset); - return offset; -} - -static void ConvertSharedMemToVgpr(IR::IREmitter& ir, IR::Inst& inst, const IR::VectorReg vgpr) { - switch (inst.GetOpcode()) { - case IR::Opcode::LoadSharedU32: - inst.ReplaceUsesWithAndRemove(ir.GetVectorReg(vgpr)); - break; - case IR::Opcode::LoadSharedU64: - inst.ReplaceUsesWithAndRemove( - ir.CompositeConstruct(ir.GetVectorReg(vgpr), ir.GetVectorReg(vgpr + 1))); - break; - case IR::Opcode::WriteSharedU32: - ir.SetVectorReg(vgpr, IR::U32{inst.Arg(1)}); - inst.Invalidate(); - break; - case IR::Opcode::WriteSharedU64: { - const auto value = inst.Arg(1); - ir.SetVectorReg(vgpr, IR::U32{ir.CompositeExtract(value, 0)}); - ir.SetVectorReg(vgpr, IR::U32{ir.CompositeExtract(value, 1)}); - inst.Invalidate(); - break; - } - default: - UNREACHABLE_MSG("Unknown shared memory opcode: {}", inst.GetOpcode()); - } -} - -void LowerSharedMemToRegisters(IR::Program& program, const RuntimeInfo& runtime_info) { - u32 next_vgpr_num = runtime_info.num_allocated_vgprs; - std::unordered_map vgpr_map; - const auto get_vgpr = [&next_vgpr_num, &vgpr_map](const u32 offset) { - const auto [it, is_new] = vgpr_map.try_emplace(offset); - if (is_new) { - ASSERT_MSG(next_vgpr_num < 256, "Out of VGPRs"); - const auto new_vgpr = static_cast(next_vgpr_num++); - it->second = new_vgpr; - } - return it->second; - }; - - for (IR::Block* const block : program.blocks) { - for (IR::Inst& inst : block->Instructions()) { - if (!IsSharedMemoryInst(inst)) { - continue; - } - const auto offset = GetSharedMemImmOffset(inst); - const auto vgpr = get_vgpr(offset); - IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)}; - ConvertSharedMemToVgpr(ir, inst, vgpr); - } - } -} - -} // namespace Shader::Optimization diff --git a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp index 029558d9e..c5bfe5796 100644 --- a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp +++ b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp @@ -78,7 +78,20 @@ bool IsDataRingInstruction(const IR::Inst& inst) { } IR::Type BufferDataType(const IR::Inst& inst, AmdGpu::NumberFormat num_format) { - return IR::Type::U32; + switch (inst.GetOpcode()) { + case IR::Opcode::LoadBufferU8: + case IR::Opcode::StoreBufferU8: + return IR::Type::U8; + case IR::Opcode::LoadBufferU16: + case IR::Opcode::StoreBufferU16: + return IR::Type::U16; + case IR::Opcode::LoadBufferFormatF32: + case IR::Opcode::StoreBufferFormatF32: + // Formatted buffer loads can use a variety of types. + return IR::Type::U32 | IR::Type::F32 | IR::Type::U16 | IR::Type::U8; + default: + return IR::Type::U32; + } } bool IsImageAtomicInstruction(const IR::Inst& inst) { @@ -121,11 +134,9 @@ public: u32 Add(const BufferResource& desc) { const u32 index{Add(buffer_resources, desc, [&desc](const auto& existing) { - // Only one GDS binding can exist. - if (desc.is_gds_buffer && existing.is_gds_buffer) { - return true; - } - return desc.sharp_idx == existing.sharp_idx && desc.inline_cbuf == existing.inline_cbuf; + return desc.sharp_idx == existing.sharp_idx && + desc.inline_cbuf == existing.inline_cbuf && + desc.buffer_type == existing.buffer_type; })}; auto& buffer = buffer_resources[index]; buffer.used_types |= desc.used_types; @@ -272,6 +283,7 @@ s32 TryHandleInlineCbuf(IR::Inst& inst, Info& info, Descriptors& descriptors, .sharp_idx = std::numeric_limits::max(), .used_types = BufferDataType(inst, cbuf.GetNumberFmt()), .inline_cbuf = cbuf, + .buffer_type = BufferType::Guest, }); } @@ -286,6 +298,7 @@ void PatchBufferSharp(IR::Block& block, IR::Inst& inst, Info& info, Descriptors& binding = descriptors.Add(BufferResource{ .sharp_idx = sharp, .used_types = BufferDataType(inst, buffer.GetNumberFmt()), + .buffer_type = BufferType::Guest, .is_written = IsBufferStore(inst), .is_formatted = inst.GetOpcode() == IR::Opcode::LoadBufferFormatF32 || inst.GetOpcode() == IR::Opcode::StoreBufferFormatF32, @@ -402,13 +415,10 @@ void PatchImageSharp(IR::Block& block, IR::Inst& inst, Info& info, Descriptors& } void PatchDataRingAccess(IR::Block& block, IR::Inst& inst, Info& info, Descriptors& descriptors) { - // Insert gds binding in the shader if it doesn't exist already. - // The buffer is used for append/consume counters. - constexpr static AmdGpu::Buffer GdsSharp{.base_address = 1}; const u32 binding = descriptors.Add(BufferResource{ .used_types = IR::Type::U32, - .inline_cbuf = GdsSharp, - .is_gds_buffer = true, + .inline_cbuf = AmdGpu::Buffer::Null(), + .buffer_type = BufferType::GdsBuffer, .is_written = true, }); @@ -420,12 +430,12 @@ void PatchDataRingAccess(IR::Block& block, IR::Inst& inst, Info& info, Descripto }; // Attempt to deduce the GDS address of counter at compile time. - const u32 gds_addr = [&] { - const IR::Value& gds_offset = inst.Arg(0); - if (gds_offset.IsImmediate()) { - // Nothing to do, offset is known. - return gds_offset.U32() & 0xFFFF; - } + u32 gds_addr = 0; + const IR::Value& gds_offset = inst.Arg(0); + if (gds_offset.IsImmediate()) { + // Nothing to do, offset is known. + gds_addr = gds_offset.U32() & 0xFFFF; + } else { const auto result = IR::BreadthFirstSearch(&inst, pred); ASSERT_MSG(result, "Unable to track M0 source"); @@ -436,8 +446,8 @@ void PatchDataRingAccess(IR::Block& block, IR::Inst& inst, Info& info, Descripto if (prod->GetOpcode() == IR::Opcode::IAdd32) { m0_val += prod->Arg(1).U32(); } - return m0_val & 0xFFFF; - }(); + gds_addr = m0_val & 0xFFFF; + } // Patch instruction. IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)}; diff --git a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp index f3a1fc9a8..219378a6c 100644 --- a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp +++ b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp @@ -74,7 +74,14 @@ void Visit(Info& info, const IR::Inst& inst) { info.uses_lane_id = true; break; case IR::Opcode::ReadConst: - info.has_readconst = true; + if (!info.has_readconst) { + info.buffers.push_back({ + .used_types = IR::Type::U32, + .inline_cbuf = AmdGpu::Buffer::Null(), + .buffer_type = BufferType::ReadConstUbo, + }); + info.has_readconst = true; + } break; case IR::Opcode::PackUfloat10_11_11: info.uses_pack_10_11_11 = true; @@ -88,10 +95,9 @@ void Visit(Info& info, const IR::Inst& inst) { } void CollectShaderInfoPass(IR::Program& program) { - Info& info{program.info}; for (IR::Block* const block : program.post_order_blocks) { for (IR::Inst& inst : block->Instructions()) { - Visit(info, inst); + Visit(program.info, inst); } } } diff --git a/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp b/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp index ec7d7e986..0ee52cf19 100644 --- a/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp +++ b/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp @@ -8,37 +8,46 @@ namespace Shader::Optimization { +static bool IsLoadShared(const IR::Inst& inst) { + return inst.GetOpcode() == IR::Opcode::LoadSharedU32 || + inst.GetOpcode() == IR::Opcode::LoadSharedU64; +} + +static bool IsWriteShared(const IR::Inst& inst) { + return inst.GetOpcode() == IR::Opcode::WriteSharedU32 || + inst.GetOpcode() == IR::Opcode::WriteSharedU64; +} + +// Inserts barriers when a shared memory write and read occur in the same basic block. static void EmitBarrierInBlock(IR::Block* block) { - // This is inteded to insert a barrier when shared memory write and read - // occur in the same basic block. Also checks if branch depth is zero as - // we don't want to insert barrier in potentially divergent code. - bool emit_barrier_on_write = false; - bool emit_barrier_on_read = false; - const auto emit_barrier = [block](bool& emit_cond, IR::Inst& inst) { - if (emit_cond) { - IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)}; - ir.Barrier(); - emit_cond = false; - } + enum class BarrierAction : u32 { + None, + BarrierOnWrite, + BarrierOnRead, }; + BarrierAction action{}; for (IR::Inst& inst : block->Instructions()) { - if (inst.GetOpcode() == IR::Opcode::LoadSharedU32 || - inst.GetOpcode() == IR::Opcode::LoadSharedU64) { - emit_barrier(emit_barrier_on_read, inst); - emit_barrier_on_write = true; + if (IsLoadShared(inst)) { + if (action == BarrierAction::BarrierOnRead) { + IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)}; + ir.Barrier(); + } + action = BarrierAction::BarrierOnWrite; + continue; } - if (inst.GetOpcode() == IR::Opcode::WriteSharedU32 || - inst.GetOpcode() == IR::Opcode::WriteSharedU64) { - emit_barrier(emit_barrier_on_write, inst); - emit_barrier_on_read = true; + if (IsWriteShared(inst)) { + if (action == BarrierAction::BarrierOnWrite) { + IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)}; + ir.Barrier(); + } + action = BarrierAction::BarrierOnRead; } } } +// Inserts a barrier after divergent conditional blocks to avoid undefined +// behavior when some threads write and others read from shared memory. static void EmitBarrierInMergeBlock(const IR::AbstractSyntaxNode::Data& data) { - // Insert a barrier after divergent conditional blocks. - // This avoids potential softlocks and crashes when some threads - // initialize shared memory and others read from it. const IR::U1 cond = data.if_node.cond; const auto insert_barrier = IR::BreadthFirstSearch(cond, [](IR::Inst* inst) -> std::optional { @@ -56,8 +65,21 @@ static void EmitBarrierInMergeBlock(const IR::AbstractSyntaxNode::Data& data) { } } -void SharedMemoryBarrierPass(IR::Program& program, const Profile& profile) { - if (!program.info.uses_shared || !profile.needs_lds_barriers) { +static constexpr u32 GcnSubgroupSize = 64; + +void SharedMemoryBarrierPass(IR::Program& program, const RuntimeInfo& runtime_info, + const Profile& profile) { + if (program.info.stage != Stage::Compute) { + return; + } + const auto& cs_info = runtime_info.cs_info; + const u32 shared_memory_size = cs_info.shared_memory_size; + const u32 threadgroup_size = + cs_info.workgroup_size[0] * cs_info.workgroup_size[1] * cs_info.workgroup_size[2]; + // The compiler can only omit barriers when the local workgroup size is the same as the HW + // subgroup. + if (shared_memory_size == 0 || threadgroup_size != GcnSubgroupSize || + !profile.needs_lds_barriers) { return; } using Type = IR::AbstractSyntaxNode::Type; @@ -67,6 +89,8 @@ void SharedMemoryBarrierPass(IR::Program& program, const Profile& profile) { --branch_depth; continue; } + // Check if branch depth is zero, we don't want to insert barrier in potentially divergent + // code. if (node.type == Type::If && branch_depth++ == 0) { EmitBarrierInMergeBlock(node.data); continue; diff --git a/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp b/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp new file mode 100644 index 000000000..25aaf257c --- /dev/null +++ b/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp @@ -0,0 +1,117 @@ +// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +#include "shader_recompiler/ir/ir_emitter.h" +#include "shader_recompiler/ir/program.h" +#include "shader_recompiler/profile.h" + +namespace Shader::Optimization { + +static bool IsSharedAccess(const IR::Inst& inst) { + const auto opcode = inst.GetOpcode(); + switch (opcode) { + case IR::Opcode::LoadSharedU32: + case IR::Opcode::LoadSharedU64: + case IR::Opcode::WriteSharedU32: + case IR::Opcode::WriteSharedU64: + case IR::Opcode::SharedAtomicAnd32: + case IR::Opcode::SharedAtomicIAdd32: + case IR::Opcode::SharedAtomicOr32: + case IR::Opcode::SharedAtomicSMax32: + case IR::Opcode::SharedAtomicUMax32: + case IR::Opcode::SharedAtomicSMin32: + case IR::Opcode::SharedAtomicUMin32: + case IR::Opcode::SharedAtomicXor32: + return true; + default: + return false; + } +} + +void SharedMemoryToStoragePass(IR::Program& program, const RuntimeInfo& runtime_info, + const Profile& profile) { + if (program.info.stage != Stage::Compute) { + return; + } + // Only perform the transform if the host shared memory is insufficient. + const u32 shared_memory_size = runtime_info.cs_info.shared_memory_size; + if (shared_memory_size <= profile.max_shared_memory_size) { + return; + } + // Add buffer binding for shared memory storage buffer. + const u32 binding = static_cast(program.info.buffers.size()); + program.info.buffers.push_back({ + .used_types = IR::Type::U32, + .inline_cbuf = AmdGpu::Buffer::Null(), + .buffer_type = BufferType::SharedMemory, + .is_written = true, + }); + for (IR::Block* const block : program.blocks) { + for (IR::Inst& inst : block->Instructions()) { + if (!IsSharedAccess(inst)) { + continue; + } + IR::IREmitter ir{*block, IR::Block::InstructionList::s_iterator_to(inst)}; + const IR::U32 handle = ir.Imm32(binding); + // Replace shared atomics first + switch (inst.GetOpcode()) { + case IR::Opcode::SharedAtomicAnd32: + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicAnd(handle, inst.Arg(0), inst.Arg(1), {})); + continue; + case IR::Opcode::SharedAtomicIAdd32: + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicIAdd(handle, inst.Arg(0), inst.Arg(1), {})); + continue; + case IR::Opcode::SharedAtomicOr32: + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicOr(handle, inst.Arg(0), inst.Arg(1), {})); + continue; + case IR::Opcode::SharedAtomicSMax32: + case IR::Opcode::SharedAtomicUMax32: { + const bool is_signed = inst.GetOpcode() == IR::Opcode::SharedAtomicSMax32; + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicIMax(handle, inst.Arg(0), inst.Arg(1), is_signed, {})); + continue; + } + case IR::Opcode::SharedAtomicSMin32: + case IR::Opcode::SharedAtomicUMin32: { + const bool is_signed = inst.GetOpcode() == IR::Opcode::SharedAtomicSMin32; + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicIMin(handle, inst.Arg(0), inst.Arg(1), is_signed, {})); + continue; + } + case IR::Opcode::SharedAtomicXor32: + inst.ReplaceUsesWithAndRemove( + ir.BufferAtomicXor(handle, inst.Arg(0), inst.Arg(1), {})); + continue; + default: + break; + } + // Replace shared operations. + const IR::U32 offset = ir.IMul(ir.GetAttributeU32(IR::Attribute::WorkgroupIndex), + ir.Imm32(shared_memory_size)); + const IR::U32 address = ir.IAdd(IR::U32{inst.Arg(0)}, offset); + switch (inst.GetOpcode()) { + case IR::Opcode::LoadSharedU32: + inst.ReplaceUsesWithAndRemove(ir.LoadBufferU32(1, handle, address, {})); + break; + case IR::Opcode::LoadSharedU64: + inst.ReplaceUsesWithAndRemove(ir.LoadBufferU32(2, handle, address, {})); + break; + case IR::Opcode::WriteSharedU32: + ir.StoreBufferU32(1, handle, address, inst.Arg(1), {}); + inst.Invalidate(); + break; + case IR::Opcode::WriteSharedU64: + ir.StoreBufferU32(2, handle, address, inst.Arg(1), {}); + inst.Invalidate(); + break; + default: + break; + } + } + } +} + +} // namespace Shader::Optimization diff --git a/src/shader_recompiler/recompiler.cpp b/src/shader_recompiler/recompiler.cpp index 5a6d1d775..1c132ebbb 100644 --- a/src/shader_recompiler/recompiler.cpp +++ b/src/shader_recompiler/recompiler.cpp @@ -65,10 +65,6 @@ IR::Program TranslateProgram(std::span code, Pools& pools, Info& info // Run optimization passes const auto stage = program.info.stage; - if (stage == Stage::Fragment) { - // Before SSA pass, as it will rewrite to VGPR load/store. - Shader::Optimization::LowerSharedMemToRegisters(program, runtime_info); - } Shader::Optimization::SsaRewritePass(program.post_order_blocks); Shader::Optimization::IdentityRemovalPass(program.blocks); if (info.l_stage == LogicalStage::TessellationControl) { @@ -90,11 +86,12 @@ IR::Program TranslateProgram(std::span code, Pools& pools, Info& info Shader::Optimization::FlattenExtendedUserdataPass(program); Shader::Optimization::ResourceTrackingPass(program); Shader::Optimization::LowerBufferFormatToRaw(program); + Shader::Optimization::SharedMemoryToStoragePass(program, runtime_info, profile); + Shader::Optimization::SharedMemoryBarrierPass(program, runtime_info, profile); Shader::Optimization::IdentityRemovalPass(program.blocks); Shader::Optimization::DeadCodeEliminationPass(program); Shader::Optimization::ConstantPropagationPass(program.post_order_blocks); Shader::Optimization::CollectShaderInfoPass(program); - Shader::Optimization::SharedMemoryBarrierPass(program, profile); return program; } diff --git a/src/shader_recompiler/runtime_info.h b/src/shader_recompiler/runtime_info.h index 78973c2d4..517392b98 100644 --- a/src/shader_recompiler/runtime_info.h +++ b/src/shader_recompiler/runtime_info.h @@ -167,6 +167,17 @@ enum class MrtSwizzle : u8 { }; static constexpr u32 MaxColorBuffers = 8; +struct PsColorBuffer { + AmdGpu::NumberFormat num_format : 4; + AmdGpu::NumberConversion num_conversion : 2; + AmdGpu::Liverpool::ShaderExportFormat export_format : 4; + u32 needs_unorm_fixup : 1; + u32 pad : 21; + AmdGpu::CompMapping swizzle; + + auto operator<=>(const PsColorBuffer&) const noexcept = default; +}; + struct FragmentRuntimeInfo { struct PsInput { u8 param_index; @@ -184,15 +195,6 @@ struct FragmentRuntimeInfo { AmdGpu::Liverpool::PsInput addr_flags; u32 num_inputs; std::array inputs; - struct PsColorBuffer { - AmdGpu::NumberFormat num_format; - AmdGpu::NumberConversion num_conversion; - AmdGpu::CompMapping swizzle; - AmdGpu::Liverpool::ShaderExportFormat export_format; - bool needs_unorm_fixup; - - auto operator<=>(const PsColorBuffer&) const noexcept = default; - }; std::array color_buffers; bool operator==(const FragmentRuntimeInfo& other) const noexcept { @@ -264,3 +266,14 @@ struct RuntimeInfo { }; } // namespace Shader + +template <> +struct fmt::formatter { + constexpr auto parse(format_parse_context& ctx) { + return ctx.begin(); + } + auto format(const Shader::Stage stage, format_context& ctx) const { + constexpr static std::array names = {"fs", "vs", "gs", "es", "hs", "ls", "cs"}; + return fmt::format_to(ctx.out(), "{}", names[static_cast(stage)]); + } +}; diff --git a/src/shader_recompiler/specialization.h b/src/shader_recompiler/specialization.h index 9bf9e71e4..1c3bfc60a 100644 --- a/src/shader_recompiler/specialization.h +++ b/src/shader_recompiler/specialization.h @@ -98,12 +98,6 @@ struct StageSpecialization { }); } u32 binding{}; - if (info->has_emulated_shared_memory) { - binding++; - } - if (info->has_readconst) { - binding++; - } ForEachSharp(binding, buffers, info->buffers, [profile_](auto& spec, const auto& desc, AmdGpu::Buffer sharp) { spec.stride = sharp.GetStride(); @@ -195,18 +189,6 @@ struct StageSpecialization { } } u32 binding{}; - if (info->has_emulated_shared_memory != other.info->has_emulated_shared_memory) { - return false; - } - if (info->has_readconst != other.info->has_readconst) { - return false; - } - if (info->has_emulated_shared_memory) { - binding++; - } - if (info->has_readconst) { - binding++; - } for (u32 i = 0; i < buffers.size(); i++) { if (other.bitset[binding++] && buffers[i] != other.buffers[i]) { return false; diff --git a/src/video_core/amdgpu/liverpool.h b/src/video_core/amdgpu/liverpool.h index 525a0c9f1..5b9b647eb 100644 --- a/src/video_core/amdgpu/liverpool.h +++ b/src/video_core/amdgpu/liverpool.h @@ -197,6 +197,10 @@ struct Liverpool { return settings.lds_dwords.Value() * 128 * 4; } + u32 NumWorkgroups() const noexcept { + return dim_x * dim_y * dim_z; + } + bool IsTgidEnabled(u32 i) const noexcept { return (settings.tgid_enable.Value() >> i) & 1; } diff --git a/src/video_core/amdgpu/resource.h b/src/video_core/amdgpu/resource.h index fa8edb3e2..64a85c812 100644 --- a/src/video_core/amdgpu/resource.h +++ b/src/video_core/amdgpu/resource.h @@ -31,6 +31,12 @@ struct Buffer { u32 _padding1 : 6; u32 type : 2; // overlaps with T# type, so should be 0 for buffer + static constexpr Buffer Null() { + Buffer buffer{}; + buffer.base_address = 1; + return buffer; + } + bool Valid() const { return type == 0u; } diff --git a/src/video_core/amdgpu/types.h b/src/video_core/amdgpu/types.h index ee2dda494..d991e0abd 100644 --- a/src/video_core/amdgpu/types.h +++ b/src/video_core/amdgpu/types.h @@ -183,7 +183,7 @@ enum class NumberFormat : u32 { Ubscaled = 13, }; -enum class CompSwizzle : u32 { +enum class CompSwizzle : u8 { Zero = 0, One = 1, Red = 4, @@ -193,10 +193,10 @@ enum class CompSwizzle : u32 { }; enum class NumberConversion : u32 { - None, - UintToUscaled, - SintToSscaled, - UnormToUbnorm, + None = 0, + UintToUscaled = 1, + SintToSscaled = 2, + UnormToUbnorm = 3, }; struct CompMapping { diff --git a/src/video_core/buffer_cache/buffer.h b/src/video_core/buffer_cache/buffer.h index ec92a0ebf..188b4b2ca 100644 --- a/src/video_core/buffer_cache/buffer.h +++ b/src/video_core/buffer_cache/buffer.h @@ -168,7 +168,7 @@ public: void Commit(); /// Maps and commits a memory region with user provided data - u64 Copy(VAddr src, size_t size, size_t alignment = 0) { + u64 Copy(auto src, size_t size, size_t alignment = 0) { const auto [data, offset] = Map(size, alignment); std::memcpy(data, reinterpret_cast(src), size); Commit(); diff --git a/src/video_core/buffer_cache/buffer_cache.cpp b/src/video_core/buffer_cache/buffer_cache.cpp index 37af62f30..ccb45c095 100644 --- a/src/video_core/buffer_cache/buffer_cache.cpp +++ b/src/video_core/buffer_cache/buffer_cache.cpp @@ -5,11 +5,8 @@ #include "common/alignment.h" #include "common/scope_exit.h" #include "common/types.h" -#include "shader_recompiler/frontend/fetch_shader.h" -#include "shader_recompiler/info.h" #include "video_core/amdgpu/liverpool.h" #include "video_core/buffer_cache/buffer_cache.h" -#include "video_core/renderer_vulkan/liverpool_to_vk.h" #include "video_core/renderer_vulkan/vk_graphics_pipeline.h" #include "video_core/renderer_vulkan/vk_instance.h" #include "video_core/renderer_vulkan/vk_scheduler.h" @@ -18,8 +15,8 @@ namespace VideoCore { static constexpr size_t DataShareBufferSize = 64_KB; -static constexpr size_t StagingBufferSize = 1_GB; -static constexpr size_t UboStreamBufferSize = 64_MB; +static constexpr size_t StagingBufferSize = 512_MB; +static constexpr size_t UboStreamBufferSize = 128_MB; BufferCache::BufferCache(const Vulkan::Instance& instance_, Vulkan::Scheduler& scheduler_, AmdGpu::Liverpool* liverpool_, TextureCache& texture_cache_, @@ -29,10 +26,8 @@ BufferCache::BufferCache(const Vulkan::Instance& instance_, Vulkan::Scheduler& s staging_buffer{instance, scheduler, MemoryUsage::Upload, StagingBufferSize}, stream_buffer{instance, scheduler, MemoryUsage::Stream, UboStreamBufferSize}, gds_buffer{instance, scheduler, MemoryUsage::Stream, 0, AllFlags, DataShareBufferSize}, - lds_buffer{instance, scheduler, MemoryUsage::DeviceLocal, 0, AllFlags, DataShareBufferSize}, memory_tracker{&tracker} { Vulkan::SetObjectName(instance.GetDevice(), gds_buffer.Handle(), "GDS Buffer"); - Vulkan::SetObjectName(instance.GetDevice(), lds_buffer.Handle(), "LDS Buffer"); // Ensure the first slot is used for the null buffer const auto null_id = @@ -251,14 +246,6 @@ void BufferCache::InlineData(VAddr address, const void* value, u32 num_bytes, bo }); } -std::pair BufferCache::ObtainHostUBO(std::span data) { - static constexpr u64 StreamThreshold = CACHING_PAGESIZE; - ASSERT(data.size_bytes() <= StreamThreshold); - const u64 offset = stream_buffer.Copy(reinterpret_cast(data.data()), data.size_bytes(), - instance.UniformMinAlignment()); - return {&stream_buffer, offset}; -} - std::pair BufferCache::ObtainBuffer(VAddr device_addr, u32 size, bool is_written, bool is_texel_buffer, BufferId buffer_id) { // For small uniform buffers that have not been modified by gpu diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index 088c22c12..71a6bed2a 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -68,9 +68,9 @@ public: return &gds_buffer; } - /// Returns a pointer to LDS device local buffer. - [[nodiscard]] const Buffer* GetLdsBuffer() const noexcept { - return &lds_buffer; + /// Retrieves the host visible device local stream buffer. + [[nodiscard]] StreamBuffer& GetStreamBuffer() noexcept { + return stream_buffer; } /// Retrieves the buffer with the specified id. @@ -90,8 +90,6 @@ public: /// Writes a value to GPU buffer. void InlineData(VAddr address, const void* value, u32 num_bytes, bool is_gds); - [[nodiscard]] std::pair ObtainHostUBO(std::span data); - /// Obtains a buffer for the specified region. [[nodiscard]] std::pair ObtainBuffer(VAddr gpu_addr, u32 size, bool is_written, bool is_texel_buffer = false, @@ -159,7 +157,6 @@ private: StreamBuffer staging_buffer; StreamBuffer stream_buffer; Buffer gds_buffer; - Buffer lds_buffer; std::shared_mutex mutex; Common::SlotVector slot_buffers; RangeSet gpu_modified_ranges; diff --git a/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp index f0346559d..f6216f54f 100644 --- a/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp +++ b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp @@ -3,11 +3,9 @@ #include -#include "video_core/buffer_cache/buffer_cache.h" #include "video_core/renderer_vulkan/vk_compute_pipeline.h" #include "video_core/renderer_vulkan/vk_instance.h" #include "video_core/renderer_vulkan/vk_scheduler.h" -#include "video_core/texture_cache/texture_cache.h" namespace Vulkan { @@ -29,23 +27,6 @@ ComputePipeline::ComputePipeline(const Instance& instance, Scheduler& scheduler, u32 binding{}; boost::container::small_vector bindings; - - if (info->has_emulated_shared_memory) { - bindings.push_back({ - .binding = binding++, - .descriptorType = vk::DescriptorType::eStorageBuffer, - .descriptorCount = 1, - .stageFlags = vk::ShaderStageFlagBits::eCompute, - }); - } - if (info->has_readconst) { - bindings.push_back({ - .binding = binding++, - .descriptorType = vk::DescriptorType::eUniformBuffer, - .descriptorCount = 1, - .stageFlags = vk::ShaderStageFlagBits::eCompute, - }); - } for (const auto& buffer : info->buffers) { const auto sharp = buffer.GetSharp(*info); bindings.push_back({ diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp index 4eecd1edf..901096259 100644 --- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp +++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp @@ -7,23 +7,27 @@ #include #include "common/assert.h" -#include "common/io_file.h" #include "shader_recompiler/backend/spirv/emit_spirv_quad_rect.h" #include "shader_recompiler/frontend/fetch_shader.h" -#include "shader_recompiler/runtime_info.h" #include "video_core/amdgpu/resource.h" -#include "video_core/buffer_cache/buffer_cache.h" #include "video_core/renderer_vulkan/vk_graphics_pipeline.h" #include "video_core/renderer_vulkan/vk_instance.h" -#include "video_core/renderer_vulkan/vk_pipeline_cache.h" #include "video_core/renderer_vulkan/vk_scheduler.h" #include "video_core/renderer_vulkan/vk_shader_util.h" -#include "video_core/texture_cache/texture_cache.h" namespace Vulkan { using Shader::Backend::SPIRV::AuxShaderType; +static constexpr std::array LogicalStageToStageBit = { + vk::ShaderStageFlagBits::eFragment, + vk::ShaderStageFlagBits::eTessellationControl, + vk::ShaderStageFlagBits::eTessellationEvaluation, + vk::ShaderStageFlagBits::eVertex, + vk::ShaderStageFlagBits::eGeometry, + vk::ShaderStageFlagBits::eCompute, +}; + GraphicsPipeline::GraphicsPipeline( const Instance& instance, Scheduler& scheduler, DescriptorHeap& desc_heap, const Shader::Profile& profile, const GraphicsPipelineKey& key_, @@ -39,7 +43,7 @@ GraphicsPipeline::GraphicsPipeline( const auto debug_str = GetDebugString(); const vk::PushConstantRange push_constants = { - .stageFlags = gp_stage_flags, + .stageFlags = AllGraphicsStageBits, .offset = 0, .size = sizeof(Shader::PushData), }; @@ -357,14 +361,7 @@ void GraphicsPipeline::BuildDescSetLayout() { if (!stage) { continue; } - if (stage->has_readconst) { - bindings.push_back({ - .binding = binding++, - .descriptorType = vk::DescriptorType::eUniformBuffer, - .descriptorCount = 1, - .stageFlags = gp_stage_flags, - }); - } + const auto stage_bit = LogicalStageToStageBit[u32(stage->l_stage)]; for (const auto& buffer : stage->buffers) { const auto sharp = buffer.GetSharp(*stage); bindings.push_back({ @@ -373,7 +370,7 @@ void GraphicsPipeline::BuildDescSetLayout() { ? vk::DescriptorType::eStorageBuffer : vk::DescriptorType::eUniformBuffer, .descriptorCount = 1, - .stageFlags = gp_stage_flags, + .stageFlags = stage_bit, }); } for (const auto& image : stage->images) { @@ -382,7 +379,7 @@ void GraphicsPipeline::BuildDescSetLayout() { .descriptorType = image.is_written ? vk::DescriptorType::eStorageImage : vk::DescriptorType::eSampledImage, .descriptorCount = 1, - .stageFlags = gp_stage_flags, + .stageFlags = stage_bit, }); } for (const auto& sampler : stage->samplers) { @@ -390,7 +387,7 @@ void GraphicsPipeline::BuildDescSetLayout() { .binding = binding++, .descriptorType = vk::DescriptorType::eSampler, .descriptorCount = 1, - .stageFlags = gp_stage_flags, + .stageFlags = stage_bit, }); } } diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.h b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h index 64cc761f4..e6596db2f 100644 --- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.h +++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h @@ -35,8 +35,7 @@ struct GraphicsPipelineKey { std::array stage_hashes; u32 num_color_attachments; std::array color_formats; - std::array - color_buffers; + std::array color_buffers; vk::Format depth_format; vk::Format stencil_format; diff --git a/src/video_core/renderer_vulkan/vk_instance.cpp b/src/video_core/renderer_vulkan/vk_instance.cpp index 780779c0b..a17f8c9c2 100644 --- a/src/video_core/renderer_vulkan/vk_instance.cpp +++ b/src/video_core/renderer_vulkan/vk_instance.cpp @@ -1,14 +1,11 @@ // SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later -#include -#include #include #include #include #include "common/assert.h" -#include "common/config.h" #include "common/debug.h" #include "sdl_window.h" #include "video_core/renderer_vulkan/liverpool_to_vk.h" @@ -206,13 +203,12 @@ std::string Instance::GetDriverVersionName() { } bool Instance::CreateDevice() { - const vk::StructureChain feature_chain = - physical_device - .getFeatures2(); + const vk::StructureChain feature_chain = physical_device.getFeatures2< + vk::PhysicalDeviceFeatures2, vk::PhysicalDeviceVulkan11Features, + vk::PhysicalDeviceVulkan12Features, vk::PhysicalDeviceRobustness2FeaturesEXT, + vk::PhysicalDeviceExtendedDynamicState3FeaturesEXT, + vk::PhysicalDevicePrimitiveTopologyListRestartFeaturesEXT, + vk::PhysicalDevicePortabilitySubsetFeaturesKHR>(); features = feature_chain.get().features; #ifdef __APPLE__ portability_features = feature_chain.get(); @@ -319,6 +315,7 @@ bool Instance::CreateDevice() { const auto topology_list_restart_features = feature_chain.get(); + const auto vk11_features = feature_chain.get(); const auto vk12_features = feature_chain.get(); vk::StructureChain device_chain = { vk::DeviceCreateInfo{ @@ -351,12 +348,17 @@ bool Instance::CreateDevice() { }, }, vk::PhysicalDeviceVulkan11Features{ - .shaderDrawParameters = true, + .storageBuffer16BitAccess = vk11_features.storageBuffer16BitAccess, + .uniformAndStorageBuffer16BitAccess = vk11_features.uniformAndStorageBuffer16BitAccess, + .shaderDrawParameters = vk11_features.shaderDrawParameters, }, vk::PhysicalDeviceVulkan12Features{ .samplerMirrorClampToEdge = vk12_features.samplerMirrorClampToEdge, .drawIndirectCount = vk12_features.drawIndirectCount, + .storageBuffer8BitAccess = vk12_features.storageBuffer8BitAccess, + .uniformAndStorageBuffer8BitAccess = vk12_features.uniformAndStorageBuffer8BitAccess, .shaderFloat16 = vk12_features.shaderFloat16, + .shaderInt8 = vk12_features.shaderInt8, .scalarBlockLayout = vk12_features.scalarBlockLayout, .uniformBufferStandardLayout = vk12_features.uniformBufferStandardLayout, .separateDepthStencilLayouts = vk12_features.separateDepthStencilLayouts, diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp index f7afd2e75..6ac7f7e43 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp @@ -345,12 +345,12 @@ bool PipelineCache::RefreshGraphicsKey() { key.color_formats[remapped_cb] = LiverpoolToVK::SurfaceFormat(col_buf.GetDataFmt(), col_buf.GetNumberFmt()); - key.color_buffers[remapped_cb] = { + key.color_buffers[remapped_cb] = Shader::PsColorBuffer{ .num_format = col_buf.GetNumberFmt(), .num_conversion = col_buf.GetNumberConversion(), - .swizzle = col_buf.Swizzle(), .export_format = regs.color_export_format.GetFormat(cb), .needs_unorm_fixup = needs_unorm_fixup, + .swizzle = col_buf.Swizzle(), }; } diff --git a/src/video_core/renderer_vulkan/vk_pipeline_common.cpp b/src/video_core/renderer_vulkan/vk_pipeline_common.cpp index bf43257f8..96e19d6a1 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_common.cpp +++ b/src/video_core/renderer_vulkan/vk_pipeline_common.cpp @@ -37,7 +37,7 @@ void Pipeline::BindResources(DescriptorWrites& set_writes, const BufferBarriers& cmdbuf.pipelineBarrier2(dependencies); } - const auto stage_flags = IsCompute() ? vk::ShaderStageFlagBits::eCompute : gp_stage_flags; + const auto stage_flags = IsCompute() ? vk::ShaderStageFlagBits::eCompute : AllGraphicsStageBits; cmdbuf.pushConstants(*pipeline_layout, stage_flags, 0u, sizeof(push_data), &push_data); // Bind descriptor set. diff --git a/src/video_core/renderer_vulkan/vk_pipeline_common.h b/src/video_core/renderer_vulkan/vk_pipeline_common.h index e9e6fed01..9633fc4ea 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_common.h +++ b/src/video_core/renderer_vulkan/vk_pipeline_common.h @@ -15,7 +15,7 @@ class BufferCache; namespace Vulkan { -static constexpr auto gp_stage_flags = +static constexpr auto AllGraphicsStageBits = vk::ShaderStageFlagBits::eVertex | vk::ShaderStageFlagBits::eTessellationControl | vk::ShaderStageFlagBits::eTessellationEvaluation | vk::ShaderStageFlagBits::eGeometry | vk::ShaderStageFlagBits::eFragment; diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index ac6aac7b3..816f149b0 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -19,6 +19,20 @@ namespace Vulkan { +static Shader::PushData MakeUserData(const AmdGpu::Liverpool::Regs& regs) { + Shader::PushData push_data{}; + push_data.step0 = regs.vgt_instance_step_rate_0; + push_data.step1 = regs.vgt_instance_step_rate_1; + + // TODO(roamic): Add support for multiple viewports and geometry shaders when ViewportIndex + // is encountered and implemented in the recompiler. + push_data.xoffset = regs.viewport_control.xoffset_enable ? regs.viewports[0].xoffset : 0.f; + push_data.xscale = regs.viewport_control.xscale_enable ? regs.viewports[0].xscale : 1.f; + push_data.yoffset = regs.viewport_control.yoffset_enable ? regs.viewports[0].yoffset : 0.f; + push_data.yscale = regs.viewport_control.yscale_enable ? regs.viewports[0].yscale : 1.f; + return push_data; +} + Rasterizer::Rasterizer(const Instance& instance_, Scheduler& scheduler_, AmdGpu::Liverpool* liverpool_) : instance{instance_}, scheduler{scheduler_}, page_manager{this}, @@ -426,95 +440,69 @@ void Rasterizer::Finish() { } bool Rasterizer::BindResources(const Pipeline* pipeline) { - buffer_infos.clear(); - buffer_views.clear(); - image_infos.clear(); - - const auto& regs = liverpool->regs; - - if (pipeline->IsCompute()) { - const auto& info = pipeline->GetStage(Shader::LogicalStage::Compute); - - // Assume if a shader reads and writes metas at the same time, it is a copy shader. - bool meta_read = false; - for (const auto& desc : info.buffers) { - if (desc.is_gds_buffer) { - continue; - } - if (!desc.is_written) { - const VAddr address = desc.GetSharp(info).base_address; - meta_read = texture_cache.IsMeta(address); - } - } - - // Most of the time when a metadata is updated with a shader it gets cleared. It means - // we can skip the whole dispatch and update the tracked state instead. Also, it is not - // intended to be consumed and in such rare cases (e.g. HTile introspection, CRAA) we - // will need its full emulation anyways. For cases of metadata read a warning will be - // logged. - if (!meta_read) { - for (const auto& desc : info.buffers) { - const auto sharp = desc.GetSharp(info); - const VAddr address = sharp.base_address; - if (desc.is_written) { - // Assume all slices were updates - if (texture_cache.ClearMeta(address)) { - LOG_TRACE(Render_Vulkan, "Metadata update skipped"); - return false; - } - } else { - if (texture_cache.IsMeta(address)) { - LOG_WARNING(Render_Vulkan, - "Unexpected metadata read by a CS shader (buffer)"); - } - } - } - } + if (IsComputeMetaClear(pipeline)) { + return false; } set_writes.clear(); buffer_barriers.clear(); + buffer_infos.clear(); + buffer_views.clear(); + image_infos.clear(); // Bind resource buffers and textures. - Shader::PushData push_data{}; Shader::Backend::Bindings binding{}; - + Shader::PushData push_data = MakeUserData(liverpool->regs); for (const auto* stage : pipeline->GetStages()) { if (!stage) { continue; } - push_data.step0 = regs.vgt_instance_step_rate_0; - push_data.step1 = regs.vgt_instance_step_rate_1; - - // TODO(roamic): add support for multiple viewports and geometry shaders when ViewportIndex - // is encountered and implemented in the recompiler. - if (stage->stage == Shader::Stage::Vertex) { - push_data.xoffset = - regs.viewport_control.xoffset_enable ? regs.viewports[0].xoffset : 0.f; - push_data.xscale = regs.viewport_control.xscale_enable ? regs.viewports[0].xscale : 1.f; - push_data.yoffset = - regs.viewport_control.yoffset_enable ? regs.viewports[0].yoffset : 0.f; - push_data.yscale = regs.viewport_control.yscale_enable ? regs.viewports[0].yscale : 1.f; - } stage->PushUd(binding, push_data); - - BindBuffers(*stage, binding, push_data, set_writes, buffer_barriers); - BindTextures(*stage, binding, set_writes); + BindBuffers(*stage, binding, push_data); + BindTextures(*stage, binding); } pipeline->BindResources(set_writes, buffer_barriers, push_data); - return true; } +bool Rasterizer::IsComputeMetaClear(const Pipeline* pipeline) { + if (!pipeline->IsCompute()) { + return false; + } + + const auto& info = pipeline->GetStage(Shader::LogicalStage::Compute); + + // Assume if a shader reads and writes metas at the same time, it is a copy shader. + for (const auto& desc : info.buffers) { + const VAddr address = desc.GetSharp(info).base_address; + if (!desc.IsSpecial() && !desc.is_written && texture_cache.IsMeta(address)) { + return false; + } + } + + // Most of the time when a metadata is updated with a shader it gets cleared. It means + // we can skip the whole dispatch and update the tracked state instead. Also, it is not + // intended to be consumed and in such rare cases (e.g. HTile introspection, CRAA) we + // will need its full emulation anyways. + for (const auto& desc : info.buffers) { + const VAddr address = desc.GetSharp(info).base_address; + if (!desc.IsSpecial() && desc.is_written && texture_cache.ClearMeta(address)) { + // Assume all slices were updates + LOG_TRACE(Render_Vulkan, "Metadata update skipped"); + return true; + } + } + return false; +} + void Rasterizer::BindBuffers(const Shader::Info& stage, Shader::Backend::Bindings& binding, - Shader::PushData& push_data, Pipeline::DescriptorWrites& set_writes, - Pipeline::BufferBarriers& buffer_barriers) { + Shader::PushData& push_data) { buffer_bindings.clear(); for (const auto& desc : stage.buffers) { const auto vsharp = desc.GetSharp(stage); - if (!desc.is_gds_buffer && vsharp.base_address != 0 && vsharp.GetSize() > 0) { + if (!desc.IsSpecial() && vsharp.base_address != 0 && vsharp.GetSize() > 0) { const auto buffer_id = buffer_cache.FindBuffer(vsharp.base_address, vsharp.GetSize()); buffer_bindings.emplace_back(buffer_id, vsharp); } else { @@ -522,47 +510,30 @@ void Rasterizer::BindBuffers(const Shader::Info& stage, Shader::Backend::Binding } } - // Bind a SSBO to act as shared memory in case of not being able to use a workgroup buffer - // (e.g. when the compute shared memory is bigger than the GPU's shared memory) - if (stage.has_emulated_shared_memory) { - const auto* lds_buf = buffer_cache.GetLdsBuffer(); - buffer_infos.emplace_back(lds_buf->Handle(), 0, lds_buf->SizeBytes()); - set_writes.push_back({ - .dstSet = VK_NULL_HANDLE, - .dstBinding = binding.unified++, - .dstArrayElement = 0, - .descriptorCount = 1, - .descriptorType = vk::DescriptorType::eStorageBuffer, - .pBufferInfo = &buffer_infos.back(), - }); - ++binding.buffer; - } - - // Bind the flattened user data buffer as a UBO so it's accessible to the shader - if (stage.has_readconst) { - const auto [vk_buffer, offset] = buffer_cache.ObtainHostUBO(stage.flattened_ud_buf); - buffer_infos.emplace_back(vk_buffer->Handle(), offset, - stage.flattened_ud_buf.size() * sizeof(u32)); - set_writes.push_back({ - .dstSet = VK_NULL_HANDLE, - .dstBinding = binding.unified++, - .dstArrayElement = 0, - .descriptorCount = 1, - .descriptorType = vk::DescriptorType::eUniformBuffer, - .pBufferInfo = &buffer_infos.back(), - }); - ++binding.buffer; - } - // Second pass to re-bind buffers that were updated after binding for (u32 i = 0; i < buffer_bindings.size(); i++) { const auto& [buffer_id, vsharp] = buffer_bindings[i]; const auto& desc = stage.buffers[i]; const bool is_storage = desc.IsStorage(vsharp, pipeline_cache.GetProfile()); + // Buffer is not from the cache, either a special buffer or unbound. if (!buffer_id) { - if (desc.is_gds_buffer) { + if (desc.buffer_type == Shader::BufferType::GdsBuffer) { const auto* gds_buf = buffer_cache.GetGdsBuffer(); buffer_infos.emplace_back(gds_buf->Handle(), 0, gds_buf->SizeBytes()); + } else if (desc.buffer_type == Shader::BufferType::ReadConstUbo) { + auto& vk_buffer = buffer_cache.GetStreamBuffer(); + const u32 ubo_size = stage.flattened_ud_buf.size() * sizeof(u32); + const u64 offset = vk_buffer.Copy(stage.flattened_ud_buf.data(), ubo_size, + instance.UniformMinAlignment()); + buffer_infos.emplace_back(vk_buffer.Handle(), offset, ubo_size); + } else if (desc.buffer_type == Shader::BufferType::SharedMemory) { + auto& lds_buffer = buffer_cache.GetStreamBuffer(); + const auto& cs_program = liverpool->GetCsRegs(); + const auto lds_size = cs_program.SharedMemSize() * cs_program.NumWorkgroups(); + const auto [data, offset] = + lds_buffer.Map(lds_size, instance.StorageMinAlignment()); + std::memset(data, 0, lds_size); + buffer_infos.emplace_back(lds_buffer.Handle(), offset, lds_size); } else if (instance.IsNullDescriptorSupported()) { buffer_infos.emplace_back(VK_NULL_HANDLE, 0, VK_WHOLE_SIZE); } else { @@ -605,8 +576,7 @@ void Rasterizer::BindBuffers(const Shader::Info& stage, Shader::Backend::Binding } } -void Rasterizer::BindTextures(const Shader::Info& stage, Shader::Backend::Bindings& binding, - Pipeline::DescriptorWrites& set_writes) { +void Rasterizer::BindTextures(const Shader::Info& stage, Shader::Backend::Bindings& binding) { image_bindings.clear(); for (const auto& image_desc : stage.images) { diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.h b/src/video_core/renderer_vulkan/vk_rasterizer.h index db458662c..292944a10 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.h +++ b/src/video_core/renderer_vulkan/vk_rasterizer.h @@ -81,11 +81,9 @@ private: bool FilterDraw(); void BindBuffers(const Shader::Info& stage, Shader::Backend::Bindings& binding, - Shader::PushData& push_data, Pipeline::DescriptorWrites& set_writes, - Pipeline::BufferBarriers& buffer_barriers); + Shader::PushData& push_data); - void BindTextures(const Shader::Info& stage, Shader::Backend::Bindings& binding, - Pipeline::DescriptorWrites& set_writes); + void BindTextures(const Shader::Info& stage, Shader::Backend::Bindings& binding); bool BindResources(const Pipeline* pipeline); void ResetBindings() { @@ -95,6 +93,8 @@ private: bound_images.clear(); } + bool IsComputeMetaClear(const Pipeline* pipeline); + private: const Instance& instance; Scheduler& scheduler;