Merge branch 'shadps4-emu:main' into shader_recompiler/format

This commit is contained in:
Daniel R. 2024-08-30 15:40:17 +02:00 committed by GitHub
commit 84f1690dfb
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
92 changed files with 2751 additions and 1848 deletions

3
.gitmodules vendored
View File

@ -82,3 +82,6 @@
path = externals/ffmpeg-core
url = https://github.com/shadps4-emu/ext-ffmpeg-core.git
shallow = true
[submodule "externals/half"]
path = externals/half
url = https://github.com/ROCm/half.git

View File

@ -521,6 +521,8 @@ set(VIDEO_CORE src/video_core/amdgpu/liverpool.cpp
src/video_core/renderer_vulkan/vk_resource_pool.h
src/video_core/renderer_vulkan/vk_scheduler.cpp
src/video_core/renderer_vulkan/vk_scheduler.h
src/video_core/renderer_vulkan/vk_shader_cache.cpp
src/video_core/renderer_vulkan/vk_shader_cache.h
src/video_core/renderer_vulkan/vk_shader_util.cpp
src/video_core/renderer_vulkan/vk_shader_util.h
src/video_core/renderer_vulkan/vk_swapchain.cpp
@ -642,6 +644,9 @@ if (APPLE)
# Replacement for std::chrono::time_zone
target_link_libraries(shadps4 PRIVATE date::date-tz)
# Half float conversions for F16C patches
target_link_libraries(shadps4 PRIVATE half)
endif()
if (NOT ENABLE_QT_GUI)

View File

@ -142,11 +142,17 @@ if (WIN32)
target_compile_options(sirit PUBLIC "-Wno-error=unused-command-line-argument")
endif()
# date
if (APPLE AND NOT TARGET date::date-tz)
if (APPLE)
# half
add_library(half INTERFACE)
target_include_directories(half INTERFACE half/include)
# date
if (NOT TARGET date::date-tz)
option(BUILD_TZ_LIB "" ON)
option(USE_SYSTEM_TZ_DB "" ON)
add_subdirectory(date)
endif()
endif()
# Tracy

1
externals/half vendored Submodule

@ -0,0 +1 @@
Subproject commit 1ddada225144cac0de8f6b5c0dd9acffd99a2e68

View File

@ -106,6 +106,7 @@ static auto UserPaths = [] {
create_path(PathType::CapturesDir, user_dir / CAPTURES_DIR);
create_path(PathType::CheatsDir, user_dir / CHEATS_DIR);
create_path(PathType::PatchesDir, user_dir / PATCHES_DIR);
create_path(PathType::AddonsDir, user_dir / ADDONS_DIR);
return paths;
}();

View File

@ -22,6 +22,7 @@ enum class PathType {
CapturesDir, // Where rdoc captures are stored.
CheatsDir, // Where cheats are stored.
PatchesDir, // Where patches are stored.
AddonsDir, // Where additional content is stored.
};
constexpr auto PORTABLE_DIR = "user";
@ -39,6 +40,7 @@ constexpr auto DOWNLOAD_DIR = "download";
constexpr auto CAPTURES_DIR = "captures";
constexpr auto CHEATS_DIR = "cheats";
constexpr auto PATCHES_DIR = "patches";
constexpr auto ADDONS_DIR = "addcont";
// Filenames
constexpr auto LOG_FILE = "shad_log.txt";

View File

@ -15,6 +15,7 @@
#else
#include <pthread.h>
#ifdef __APPLE__
#include <half.hpp>
#include <sys/sysctl.h>
#endif
#endif
@ -30,6 +31,12 @@ static Xbyak::Reg ZydisToXbyakRegister(const ZydisRegister reg) {
if (reg >= ZYDIS_REGISTER_RAX && reg <= ZYDIS_REGISTER_R15) {
return Xbyak::Reg64(reg - ZYDIS_REGISTER_RAX + Xbyak::Operand::RAX);
}
if (reg >= ZYDIS_REGISTER_XMM0 && reg <= ZYDIS_REGISTER_XMM31) {
return Xbyak::Xmm(reg - ZYDIS_REGISTER_XMM0 + xmm0.getIdx());
}
if (reg >= ZYDIS_REGISTER_YMM0 && reg <= ZYDIS_REGISTER_YMM31) {
return Xbyak::Ymm(reg - ZYDIS_REGISTER_YMM0 + ymm0.getIdx());
}
UNREACHABLE_MSG("Unsupported register: {}", static_cast<u32>(reg));
}
@ -66,6 +73,12 @@ static Xbyak::Address ZydisToXbyakMemoryOperand(const ZydisDecodedOperand& opera
return ptr[expression];
}
static u64 ZydisToXbyakImmediateOperand(const ZydisDecodedOperand& operand) {
ASSERT_MSG(operand.type == ZYDIS_OPERAND_TYPE_IMMEDIATE,
"Expected immediate operand, got type: {}", static_cast<u32>(operand.type));
return operand.imm.value.u;
}
static std::unique_ptr<Xbyak::Operand> ZydisToXbyakOperand(const ZydisDecodedOperand& operand) {
switch (operand.type) {
case ZYDIS_OPERAND_TYPE_REGISTER: {
@ -110,51 +123,135 @@ static Xbyak::Reg AllocateScratchRegister(
#ifdef __APPLE__
static constexpr u32 MaxSavedRegisters = 3;
static pthread_key_t register_save_slots[MaxSavedRegisters];
static std::once_flag register_save_init_flag;
static pthread_key_t stack_pointer_slot;
static pthread_key_t patch_stack_slot;
static std::once_flag patch_context_slots_init_flag;
static_assert(sizeof(void*) == sizeof(u64),
"Cannot fit a register inside a thread local storage slot.");
static void InitializeRegisterSaveSlots() {
for (u32 i = 0; i < MaxSavedRegisters; i++) {
ASSERT_MSG(pthread_key_create(&register_save_slots[i], nullptr) == 0,
"Unable to allocate thread-local register save slot {}", i);
static void InitializePatchContextSlots() {
ASSERT_MSG(pthread_key_create(&stack_pointer_slot, nullptr) == 0,
"Unable to allocate thread-local register for stack pointer.");
ASSERT_MSG(pthread_key_create(&patch_stack_slot, nullptr) == 0,
"Unable to allocate thread-local register for patch stack.");
}
void InitializeThreadPatchStack() {
std::call_once(patch_context_slots_init_flag, InitializePatchContextSlots);
const auto* patch_stack = std::malloc(0x1000);
pthread_setspecific(patch_stack_slot, patch_stack);
}
void CleanupThreadPatchStack() {
std::call_once(patch_context_slots_init_flag, InitializePatchContextSlots);
auto* patch_stack = pthread_getspecific(patch_stack_slot);
if (patch_stack != nullptr) {
std::free(patch_stack);
pthread_setspecific(patch_stack_slot, nullptr);
}
}
static void SaveRegisters(Xbyak::CodeGenerator& c, const std::initializer_list<Xbyak::Reg> regs) {
ASSERT_MSG(regs.size() <= MaxSavedRegisters, "Not enough space to save {} registers.",
regs.size());
std::call_once(register_save_init_flag, &InitializeRegisterSaveSlots);
u32 index = 0;
for (const auto& reg : regs) {
const auto offset = reinterpret_cast<void*>(register_save_slots[index++] * sizeof(void*));
/// Saves the stack pointer to thread local storage and loads the patch stack.
static void SaveStack(Xbyak::CodeGenerator& c) {
std::call_once(patch_context_slots_init_flag, InitializePatchContextSlots);
// Save stack pointer and load patch stack.
c.putSeg(gs);
c.mov(qword[offset], reg.cvt64());
}
c.mov(qword[reinterpret_cast<void*>(stack_pointer_slot * sizeof(void*))], rsp);
c.putSeg(gs);
c.mov(rsp, qword[reinterpret_cast<void*>(patch_stack_slot * sizeof(void*))]);
}
/// Restores the stack pointer from thread local storage.
static void RestoreStack(Xbyak::CodeGenerator& c) {
std::call_once(patch_context_slots_init_flag, InitializePatchContextSlots);
// Save patch stack pointer and load original stack.
c.putSeg(gs);
c.mov(qword[reinterpret_cast<void*>(patch_stack_slot * sizeof(void*))], rsp);
c.putSeg(gs);
c.mov(rsp, qword[reinterpret_cast<void*>(stack_pointer_slot * sizeof(void*))]);
}
#else
// These utilities are not implemented as we can't save anything to thread local storage without
// temporary registers.
void InitializeThreadPatchStack() {
// No-op
}
void CleanupThreadPatchStack() {
// No-op
}
/// Saves the stack pointer to thread local storage and loads the patch stack.
static void SaveStack(Xbyak::CodeGenerator& c) {
UNIMPLEMENTED();
}
/// Restores the stack pointer from thread local storage.
static void RestoreStack(Xbyak::CodeGenerator& c) {
UNIMPLEMENTED();
}
#endif
/// Switches to the patch stack, saves registers, and restores the original stack.
static void SaveRegisters(Xbyak::CodeGenerator& c, const std::initializer_list<Xbyak::Reg> regs) {
SaveStack(c);
for (const auto& reg : regs) {
c.push(reg.cvt64());
}
RestoreStack(c);
}
/// Switches to the patch stack, restores registers, and restores the original stack.
static void RestoreRegisters(Xbyak::CodeGenerator& c,
const std::initializer_list<Xbyak::Reg> regs) {
ASSERT_MSG(regs.size() <= MaxSavedRegisters, "Not enough space to restore {} registers.",
regs.size());
std::call_once(register_save_init_flag, &InitializeRegisterSaveSlots);
u32 index = 0;
SaveStack(c);
for (const auto& reg : regs) {
const auto offset = reinterpret_cast<void*>(register_save_slots[index++] * sizeof(void*));
c.pop(reg.cvt64());
}
RestoreStack(c);
}
c.putSeg(gs);
c.mov(reg.cvt64(), qword[offset]);
/// Switches to the patch stack and stores all registers.
static void SaveContext(Xbyak::CodeGenerator& c) {
SaveStack(c);
for (int reg = Xbyak::Operand::RAX; reg <= Xbyak::Operand::R15; reg++) {
c.push(Xbyak::Reg64(reg));
}
for (int reg = 0; reg <= 7; reg++) {
c.sub(rsp, 32);
c.vmovdqu(ptr[rsp], Xbyak::Ymm(reg));
}
}
/// Restores all registers and restores the original stack.
/// If the destination is a register, it is not restored to preserve the output.
static void RestoreContext(Xbyak::CodeGenerator& c, const Xbyak::Operand& dst) {
for (int reg = 7; reg >= 0; reg--) {
if ((!dst.isXMM() && !dst.isYMM()) || dst.getIdx() != reg) {
c.vmovdqu(Xbyak::Ymm(reg), ptr[rsp]);
}
c.add(rsp, 32);
}
for (int reg = Xbyak::Operand::R15; reg >= Xbyak::Operand::RAX; reg--) {
if (!dst.isREG() || dst.getIdx() != reg) {
c.pop(Xbyak::Reg64(reg));
} else {
c.add(rsp, 4);
}
}
RestoreStack(c);
}
#ifdef __APPLE__
static void GenerateANDN(const ZydisDecodedOperand* operands, Xbyak::CodeGenerator& c) {
const auto dst = ZydisToXbyakRegisterOperand(operands[0]);
const auto src1 = ZydisToXbyakRegisterOperand(operands[1]);
@ -204,9 +301,9 @@ static void GenerateBEXTR(const ZydisDecodedOperand* operands, Xbyak::CodeGenera
c.and_(dst, scratch2);
if (dst.getIdx() == shift.getIdx()) {
RestoreRegisters(c, {scratch1, scratch2});
RestoreRegisters(c, {scratch2, scratch1});
} else {
RestoreRegisters(c, {scratch1, scratch2, shift});
RestoreRegisters(c, {shift, scratch2, scratch1});
}
}
@ -258,10 +355,138 @@ static void GenerateBLSR(const ZydisDecodedOperand* operands, Xbyak::CodeGenerat
RestoreRegisters(c, {scratch});
}
bool FilterRosetta2Only(const ZydisDecodedOperand*) {
static __attribute__((sysv_abi)) void PerformVCVTPH2PS(float* out, const half_float::half* in,
const u32 count) {
for (u32 i = 0; i < count; i++) {
out[i] = half_float::half_cast<float>(in[i]);
}
}
static void GenerateVCVTPH2PS(const ZydisDecodedOperand* operands, Xbyak::CodeGenerator& c) {
const auto dst = ZydisToXbyakRegisterOperand(operands[0]);
const auto src = ZydisToXbyakOperand(operands[1]);
const auto float_count = dst.getBit() / 32;
const auto byte_count = float_count * 4;
SaveContext(c);
// Allocate stack space for outputs and load into first parameter.
c.sub(rsp, byte_count);
c.mov(rdi, rsp);
if (src->isXMM()) {
// Allocate stack space for inputs and load into second parameter.
c.sub(rsp, byte_count);
c.mov(rsi, rsp);
// Move input to the allocated space.
c.movdqu(ptr[rsp], *reinterpret_cast<Xbyak::Xmm*>(src.get()));
} else {
c.lea(rsi, src->getAddress());
}
// Load float count into third parameter.
c.mov(rdx, float_count);
c.mov(rax, reinterpret_cast<u64>(PerformVCVTPH2PS));
c.call(rax);
if (src->isXMM()) {
// Clean up after inputs space.
c.add(rsp, byte_count);
}
// Load outputs into destination register and clean up space.
if (dst.isYMM()) {
c.vmovdqu(*reinterpret_cast<const Xbyak::Ymm*>(&dst), ptr[rsp]);
} else {
c.movdqu(*reinterpret_cast<const Xbyak::Xmm*>(&dst), ptr[rsp]);
}
c.add(rsp, byte_count);
RestoreContext(c, dst);
}
using SingleToHalfFloatConverter = half_float::half (*)(float);
static const SingleToHalfFloatConverter SingleToHalfFloatConverters[4] = {
half_float::half_cast<half_float::half, std::round_to_nearest, float>,
half_float::half_cast<half_float::half, std::round_toward_neg_infinity, float>,
half_float::half_cast<half_float::half, std::round_toward_infinity, float>,
half_float::half_cast<half_float::half, std::round_toward_zero, float>,
};
static __attribute__((sysv_abi)) void PerformVCVTPS2PH(half_float::half* out, const float* in,
const u32 count, const u8 rounding_mode) {
const auto conversion_func = SingleToHalfFloatConverters[rounding_mode];
for (u32 i = 0; i < count; i++) {
out[i] = conversion_func(in[i]);
}
}
static void GenerateVCVTPS2PH(const ZydisDecodedOperand* operands, Xbyak::CodeGenerator& c) {
const auto dst = ZydisToXbyakOperand(operands[0]);
const auto src = ZydisToXbyakRegisterOperand(operands[1]);
const auto ctrl = ZydisToXbyakImmediateOperand(operands[2]);
const auto float_count = src.getBit() / 32;
const auto byte_count = float_count * 4;
SaveContext(c);
if (dst->isXMM()) {
// Allocate stack space for outputs and load into first parameter.
c.sub(rsp, byte_count);
c.mov(rdi, rsp);
} else {
c.lea(rdi, dst->getAddress());
}
// Allocate stack space for inputs and load into second parameter.
c.sub(rsp, byte_count);
c.mov(rsi, rsp);
// Move input to the allocated space.
if (src.isYMM()) {
c.vmovdqu(ptr[rsp], *reinterpret_cast<const Xbyak::Ymm*>(&src));
} else {
c.movdqu(ptr[rsp], *reinterpret_cast<const Xbyak::Xmm*>(&src));
}
// Load float count into third parameter.
c.mov(rdx, float_count);
// Load rounding mode into fourth parameter.
if (ctrl & 4) {
// Load from MXCSR.RC.
c.stmxcsr(ptr[rsp - 4]);
c.mov(rcx, ptr[rsp - 4]);
c.shr(rcx, 13);
c.and_(rcx, 3);
} else {
c.mov(rcx, ctrl & 3);
}
c.mov(rax, reinterpret_cast<u64>(PerformVCVTPS2PH));
c.call(rax);
// Clean up after inputs space.
c.add(rsp, byte_count);
if (dst->isXMM()) {
// Load outputs into destination register and clean up space.
c.movdqu(*reinterpret_cast<Xbyak::Xmm*>(dst.get()), ptr[rsp]);
c.add(rsp, byte_count);
}
RestoreContext(c, *dst);
}
static bool FilterRosetta2Only(const ZydisDecodedOperand*) {
int ret = 0;
size_t size = sizeof(ret);
if (sysctlbyname("sysctl.proc_translated", &ret, &size, NULL, 0) != 0) {
if (sysctlbyname("sysctl.proc_translated", &ret, &size, nullptr, 0) != 0) {
return false;
}
return ret;
@ -339,12 +564,16 @@ static const std::unordered_map<ZydisMnemonic, PatchInfo> Patches = {
#endif
#ifdef __APPLE__
// BMI1 instructions that are not supported by Rosetta 2 on Apple Silicon.
// Patches for instruction sets not supported by Rosetta 2.
// BMI1
{ZYDIS_MNEMONIC_ANDN, {FilterRosetta2Only, GenerateANDN, true}},
{ZYDIS_MNEMONIC_BEXTR, {FilterRosetta2Only, GenerateBEXTR, true}},
{ZYDIS_MNEMONIC_BLSI, {FilterRosetta2Only, GenerateBLSI, true}},
{ZYDIS_MNEMONIC_BLSMSK, {FilterRosetta2Only, GenerateBLSMSK, true}},
{ZYDIS_MNEMONIC_BLSR, {FilterRosetta2Only, GenerateBLSR, true}},
// F16C
{ZYDIS_MNEMONIC_VCVTPH2PS, {FilterRosetta2Only, GenerateVCVTPH2PS, true}},
{ZYDIS_MNEMONIC_VCVTPS2PH, {FilterRosetta2Only, GenerateVCVTPS2PH, true}},
#endif
};

View File

@ -9,6 +9,12 @@ class CodeGenerator;
namespace Core {
/// Initializes a stack for the current thread for use by patch implementations.
void InitializeThreadPatchStack();
/// Cleans up the patch stack for the current thread.
void CleanupThreadPatchStack();
/// Patches CPU instructions that cannot run as-is on the host.
void PatchInstructions(u64 segment_addr, u64 segment_size, Xbyak::CodeGenerator& c);

View File

@ -67,15 +67,19 @@ bool PKG::Open(const std::filesystem::path& filepath) {
file.Seek(0x47); // skip first 7 characters of content_id
file.Read(pkgTitleID);
file.Seek(0);
pkg.resize(pkgheader.pkg_promote_size);
file.Read(pkg);
u32 offset = pkgheader.pkg_table_entry_offset;
u32 n_files = pkgheader.pkg_table_entry_count;
file.Seek(offset);
for (int i = 0; i < n_files; i++) {
PKGEntry entry;
std::memcpy(&entry, &pkg[offset + i * 0x20], sizeof(entry));
PKGEntry entry{};
file.Read(entry.id);
file.Read(entry.filename_offset);
file.Read(entry.flags1);
file.Read(entry.flags2);
file.Read(entry.offset);
file.Read(entry.size);
file.Seek(8, Common::FS::SeekOrigin::CurrentPosition);
// Try to figure out the name
const auto name = GetEntryNameByType(entry.id);
@ -113,9 +117,6 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
failreason = "Content size is bigger than pkg size";
return false;
}
file.Seek(0);
pkg.resize(pkgheader.pkg_promote_size);
file.Read(pkg);
u32 offset = pkgheader.pkg_table_entry_offset;
u32 n_files = pkgheader.pkg_table_entry_count;
@ -126,9 +127,18 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
std::array<std::array<u8, 256>, 7> key1;
std::array<u8, 256> imgkeydata;
file.Seek(offset);
for (int i = 0; i < n_files; i++) {
PKGEntry entry;
std::memcpy(&entry, &pkg[offset + i * 0x20], sizeof(entry));
PKGEntry entry{};
file.Read(entry.id);
file.Read(entry.filename_offset);
file.Read(entry.flags1);
file.Read(entry.flags2);
file.Read(entry.offset);
file.Read(entry.size);
file.Seek(8, Common::FS::SeekOrigin::CurrentPosition);
auto currentPos = file.Tell();
// Try to figure out the name
const auto name = GetEntryNameByType(entry.id);
@ -139,8 +149,15 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
// Just print with id
Common::FS::IOFile out(extract_path / "sce_sys" / std::to_string(entry.id),
Common::FS::FileAccessMode::Write);
out.WriteRaw<u8>(pkg.data() + entry.offset, entry.size);
file.Seek(entry.offset);
std::vector<u8> data;
data.resize(entry.size);
file.ReadRaw<u8>(data.data(), entry.size);
out.WriteRaw<u8>(data.data(), entry.size);
out.Close();
file.Seek(currentPos);
continue;
}
@ -178,14 +195,25 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
}
Common::FS::IOFile out(extract_path / "sce_sys" / name, Common::FS::FileAccessMode::Write);
out.WriteRaw<u8>(pkg.data() + entry.offset, entry.size);
file.Seek(entry.offset);
std::vector<u8> data;
data.resize(entry.size);
file.ReadRaw<u8>(data.data(), entry.size);
out.WriteRaw<u8>(data.data(), entry.size);
out.Close();
// Decrypt Np stuff and overwrite.
if (entry.id == 0x400 || entry.id == 0x401 || entry.id == 0x402 ||
entry.id == 0x403) { // somehow 0x401 is not decrypting
decNp.resize(entry.size);
std::span<u8> cipherNp(pkg.data() + entry.offset, entry.size);
file.Seek(entry.offset);
std::vector<u8> data;
data.resize(entry.size);
file.ReadRaw<u8>(data.data(), entry.size);
std::span<u8> cipherNp(data.data(), entry.size);
std::array<u8, 64> concatenated_ivkey_dk3_;
std::memcpy(concatenated_ivkey_dk3_.data(), &entry, sizeof(entry));
std::memcpy(concatenated_ivkey_dk3_.data() + sizeof(entry), dk3_.data(), sizeof(dk3_));
@ -197,6 +225,8 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
out.Write(decNp);
out.Close();
}
file.Seek(currentPos);
}
// Extract trophy files
@ -214,6 +244,9 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
PKG::crypto.PfsGenCryptoKey(ekpfsKey, seed, dataKey, tweakKey);
const u32 length = pkgheader.pfs_cache_size * 0x2; // Seems to be ok.
int num_blocks = 0;
std::vector<u8> pfsc(length);
if (length != 0) {
// Read encrypted pfs_image
std::vector<u8> pfs_encrypted(length);
file.Seek(pkgheader.pfs_image_offset);
@ -225,18 +258,18 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
// Retrieve PFSC from decrypted pfs_image.
pfsc_offset = GetPFSCOffset(pfs_decrypted);
std::vector<u8> pfsc(length);
std::memcpy(pfsc.data(), pfs_decrypted.data() + pfsc_offset, length - pfsc_offset);
PFSCHdr pfsChdr;
std::memcpy(&pfsChdr, pfsc.data(), sizeof(pfsChdr));
const int num_blocks = (int)(pfsChdr.data_length / pfsChdr.block_sz2);
num_blocks = (int)(pfsChdr.data_length / pfsChdr.block_sz2);
sectorMap.resize(num_blocks + 1); // 8 bytes, need extra 1 to get the last offset.
for (int i = 0; i < num_blocks + 1; i++) {
std::memcpy(&sectorMap[i], pfsc.data() + pfsChdr.block_offsets + i * 8, 8);
}
}
u32 ent_size = 0;
u32 ndinode = 0;
@ -296,7 +329,15 @@ bool PKG::Extract(const std::filesystem::path& filepath, const std::filesystem::
} else {
// Set the the folder according to the current inode.
// Can be 2 or more (rarely)
extractPaths[ndinode_counter] = extract_path.parent_path() / GetTitleID();
auto parent_path = extract_path.parent_path();
auto title_id = GetTitleID();
if (parent_path.filename() != title_id) {
extractPaths[ndinode_counter] = parent_path / title_id;
} else {
// DLCs path has different structure
extractPaths[ndinode_counter] = extract_path;
}
uroot_reached = false;
break;
}

View File

@ -149,7 +149,6 @@ public:
private:
Crypto crypto;
TRP trp;
std::vector<u8> pkg;
u64 pkgSize = 0;
char pkgTitleID[9];
PKGHeader pkgheader;

View File

@ -7,14 +7,33 @@
#include <common/singleton.h>
#include <core/file_format/psf.h>
#include <core/file_sys/fs.h>
#include "app_content.h"
#include "common/io_file.h"
#include "common/logging/log.h"
#include "common/string_util.h"
#include "core/libraries/error_codes.h"
#include "core/libraries/libs.h"
namespace Libraries::AppContent {
int32_t addcont_count = 0;
struct AddContInfo {
char entitlementLabel[ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE];
OrbisAppContentAddcontDownloadStatus status;
OrbisAppContentGetEntitlementKey key;
};
std::array<AddContInfo, ORBIS_APP_CONTENT_INFO_LIST_MAX_SIZE> addcont_info = {{
{"0000000000000000",
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_INSTALLED,
{0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00}},
}};
std::string title_id;
int PS4_SYSV_ABI _Z5dummyv() {
LOG_ERROR(Lib_AppContent, "(STUBBED) called");
return ORBIS_OK;
@ -35,9 +54,31 @@ int PS4_SYSV_ABI sceAppContentAddcontEnqueueDownloadSp() {
return ORBIS_OK;
}
int PS4_SYSV_ABI sceAppContentAddcontMount() {
LOG_ERROR(Lib_AppContent, "(STUBBED) called");
int PS4_SYSV_ABI sceAppContentAddcontMount(u32 service_label,
const OrbisNpUnifiedEntitlementLabel* entitlement_label,
OrbisAppContentMountPoint* mount_point) {
LOG_INFO(Lib_AppContent, "called");
const auto& mount_dir = Common::FS::GetUserPath(Common::FS::PathType::AddonsDir) / title_id /
entitlement_label->data;
auto* mnt = Common::Singleton<Core::FileSys::MntPoints>::Instance();
for (int i = 0; i < addcont_count; i++) {
if (strncmp(entitlement_label->data, addcont_info[i].entitlementLabel,
ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE - 1) != 0) {
continue;
}
if (addcont_info[i].status != ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_INSTALLED) {
return ORBIS_APP_CONTENT_ERROR_NOT_FOUND;
}
snprintf(mount_point->data, ORBIS_APP_CONTENT_MOUNTPOINT_DATA_MAXSIZE, "/addcont%d", i);
mnt->Mount(mount_dir, mount_point->data);
return ORBIS_OK;
}
return ORBIS_APP_CONTENT_ERROR_NOT_FOUND;
}
int PS4_SYSV_ABI sceAppContentAddcontShrink() {
@ -124,22 +165,80 @@ int PS4_SYSV_ABI sceAppContentGetAddcontDownloadProgress() {
return ORBIS_OK;
}
int PS4_SYSV_ABI sceAppContentGetAddcontInfo() {
LOG_ERROR(Lib_AppContent, "(STUBBED) called");
int PS4_SYSV_ABI sceAppContentGetAddcontInfo(u32 service_label,
const OrbisNpUnifiedEntitlementLabel* entitlementLabel,
OrbisAppContentAddcontInfo* info) {
LOG_INFO(Lib_AppContent, "called");
if (entitlementLabel == nullptr || info == nullptr) {
return ORBIS_APP_CONTENT_ERROR_PARAMETER;
}
for (auto i = 0; i < addcont_count; i++) {
if (strncmp(entitlementLabel->data, addcont_info[i].entitlementLabel,
ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE - 1) != 0) {
continue;
}
LOG_INFO(Lib_AppContent, "found DLC {}", entitlementLabel->data);
strncpy(info->entitlement_label.data, addcont_info[i].entitlementLabel,
ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE);
info->status = addcont_info[i].status;
return ORBIS_OK;
}
return ORBIS_APP_CONTENT_ERROR_DRM_NO_ENTITLEMENT;
}
int PS4_SYSV_ABI sceAppContentGetAddcontInfoList(u32 service_label,
OrbisAppContentAddcontInfo* list, u32 list_num,
u32* hit_num) {
*hit_num = 0;
LOG_ERROR(Lib_AppContent, "(DUMMY) called");
LOG_INFO(Lib_AppContent, "called");
if (list_num == 0 || list == nullptr) {
if (hit_num == nullptr) {
return ORBIS_APP_CONTENT_ERROR_PARAMETER;
}
*hit_num = addcont_count;
return ORBIS_OK;
}
int dlcs_to_list = addcont_count < list_num ? addcont_count : list_num;
for (int i = 0; i < dlcs_to_list; i++) {
strncpy(list[i].entitlement_label.data, addcont_info[i].entitlementLabel,
ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE);
list[i].status = addcont_info[i].status;
}
if (hit_num != nullptr) {
*hit_num = dlcs_to_list;
}
return ORBIS_OK;
}
int PS4_SYSV_ABI sceAppContentGetEntitlementKey() {
LOG_ERROR(Lib_AppContent, "(STUBBED) called");
int PS4_SYSV_ABI sceAppContentGetEntitlementKey(
u32 service_label, const OrbisNpUnifiedEntitlementLabel* entitlement_label,
OrbisAppContentGetEntitlementKey* key) {
LOG_ERROR(Lib_AppContent, "called");
if (entitlement_label == nullptr || key == nullptr) {
return ORBIS_APP_CONTENT_ERROR_PARAMETER;
}
for (int i = 0; i < addcont_count; i++) {
if (strncmp(entitlement_label->data, addcont_info[i].entitlementLabel,
ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE - 1) != 0) {
continue;
}
memcpy(key->data, addcont_info[i].key.data, ORBIS_APP_CONTENT_ENTITLEMENT_KEY_SIZE);
return ORBIS_OK;
}
return ORBIS_APP_CONTENT_ERROR_DRM_NO_ENTITLEMENT;
}
int PS4_SYSV_ABI sceAppContentGetRegion() {
@ -150,7 +249,25 @@ int PS4_SYSV_ABI sceAppContentGetRegion() {
int PS4_SYSV_ABI sceAppContentInitialize(const OrbisAppContentInitParam* initParam,
OrbisAppContentBootParam* bootParam) {
LOG_ERROR(Lib_AppContent, "(DUMMY) called");
bootParam->attr = 0; // always 0
auto* param_sfo = Common::Singleton<PSF>::Instance();
const auto addons_dir = Common::FS::GetUserPath(Common::FS::PathType::AddonsDir);
title_id = param_sfo->GetString("TITLE_ID");
auto addon_path = addons_dir / title_id;
if (std::filesystem::exists(addon_path)) {
for (const auto& entry : std::filesystem::directory_iterator(addon_path)) {
if (entry.is_directory()) {
auto entitlement_label = entry.path().filename().string();
AddContInfo info{};
info.status = ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_INSTALLED;
strcpy(info.entitlementLabel, entitlement_label.c_str());
addcont_info[addcont_count++] = info;
}
}
}
return ORBIS_OK;
}

View File

@ -41,6 +41,16 @@ struct OrbisAppContentMountPoint {
constexpr int ORBIS_APP_CONTENT_TEMPORARY_DATA_OPTION_NONE = 0;
constexpr int ORBIS_APP_CONTENT_TEMPORARY_DATA_OPTION_FORMAT = (1 << 0);
constexpr int ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE = 17;
constexpr int ORBIS_APP_CONTENT_ENTITLEMENT_KEY_SIZE = 16;
constexpr int ORBIS_APP_CONTENT_INFO_LIST_MAX_SIZE = 2500;
enum OrbisAppContentAddcontDownloadStatus : u32 {
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_NO_EXTRA_DATA = 0,
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_NO_IN_QUEUE = 1,
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_DOWNLOADING = 2,
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_DOWNLOAD_SUSPENDED = 3,
ORBIS_APP_CONTENT_ADDCONT_DOWNLOAD_STATUS_INSTALLED = 4
};
struct OrbisNpUnifiedEntitlementLabel {
char data[ORBIS_NP_UNIFIED_ENTITLEMENT_LABEL_SIZE];
@ -54,11 +64,17 @@ struct OrbisAppContentAddcontInfo {
u32 status;
};
struct OrbisAppContentGetEntitlementKey {
char data[ORBIS_APP_CONTENT_ENTITLEMENT_KEY_SIZE];
};
int PS4_SYSV_ABI _Z5dummyv();
int PS4_SYSV_ABI sceAppContentAddcontDelete();
int PS4_SYSV_ABI sceAppContentAddcontEnqueueDownload();
int PS4_SYSV_ABI sceAppContentAddcontEnqueueDownloadSp();
int PS4_SYSV_ABI sceAppContentAddcontMount();
int PS4_SYSV_ABI sceAppContentAddcontMount(u32 service_label,
const OrbisNpUnifiedEntitlementLabel* entitlement_label,
OrbisAppContentMountPoint* mount_point);
int PS4_SYSV_ABI sceAppContentAddcontShrink();
int PS4_SYSV_ABI sceAppContentAddcontUnmount();
int PS4_SYSV_ABI sceAppContentAppParamGetInt(OrbisAppContentAppParamId paramId, s32* value);
@ -70,11 +86,15 @@ int PS4_SYSV_ABI sceAppContentDownload1Shrink();
int PS4_SYSV_ABI sceAppContentDownloadDataFormat();
int PS4_SYSV_ABI sceAppContentDownloadDataGetAvailableSpaceKb();
int PS4_SYSV_ABI sceAppContentGetAddcontDownloadProgress();
int PS4_SYSV_ABI sceAppContentGetAddcontInfo();
int PS4_SYSV_ABI sceAppContentGetAddcontInfo(u32 service_label,
const OrbisNpUnifiedEntitlementLabel* entitlementLabel,
OrbisAppContentAddcontInfo* info);
int PS4_SYSV_ABI sceAppContentGetAddcontInfoList(u32 service_label,
OrbisAppContentAddcontInfo* list, u32 list_num,
u32* hit_num);
int PS4_SYSV_ABI sceAppContentGetEntitlementKey();
int PS4_SYSV_ABI sceAppContentGetEntitlementKey(
u32 service_label, const OrbisNpUnifiedEntitlementLabel* entitlement_label,
OrbisAppContentGetEntitlementKey* key);
int PS4_SYSV_ABI sceAppContentGetRegion();
int PS4_SYSV_ABI sceAppContentInitialize(const OrbisAppContentInitParam* initParam,
OrbisAppContentBootParam* bootParam);

View File

@ -461,3 +461,5 @@ constexpr int ORBIS_AVPLAYER_ERROR_INFO_OTHER_ENCRY = 0x806A00BF;
// AppContent library
constexpr int ORBIS_APP_CONTENT_ERROR_PARAMETER = 0x80D90002;
constexpr int ORBIS_APP_CONTENT_ERROR_DRM_NO_ENTITLEMENT = 0x80D90007;
constexpr int ORBIS_APP_CONTENT_ERROR_NOT_FOUND = 0x80D90005;

View File

@ -10,6 +10,7 @@
#include "common/logging/log.h"
#include "common/singleton.h"
#include "common/thread.h"
#include "core/cpu_patches.h"
#include "core/libraries/error_codes.h"
#include "core/libraries/kernel/libkernel.h"
#include "core/libraries/kernel/thread_management.h"
@ -985,6 +986,7 @@ static void cleanup_thread(void* arg) {
destructor(value);
}
}
Core::CleanupThreadPatchStack();
thread->is_almost_done = true;
}
@ -992,6 +994,7 @@ static void* run_thread(void* arg) {
auto* thread = static_cast<ScePthread>(arg);
Common::SetCurrentThreadName(thread->name.c_str());
auto* linker = Common::Singleton<Core::Linker>::Instance();
Core::InitializeThreadPatchStack();
linker->InitTlsForThread(false);
void* ret = nullptr;
g_pthread_self = thread;

View File

@ -2,9 +2,8 @@
// SPDX-License-Identifier: GPL-2.0-or-later
#include <condition_variable>
#include <list>
#include <mutex>
#include <utility>
#include <boost/intrusive/list.hpp>
#include <pthread.h>
#include "common/assert.h"
#include "common/logging/log.h"
@ -13,9 +12,6 @@
namespace Libraries::Kernel {
using ListBaseHook =
boost::intrusive::list_base_hook<boost::intrusive::link_mode<boost::intrusive::normal_link>>;
class Semaphore {
public:
Semaphore(s32 init_count, s32 max_count, std::string_view name, bool is_fifo)
@ -37,7 +33,7 @@ public:
// Create waiting thread object and add it into the list of waiters.
WaitingThread waiter{need_count, is_fifo};
AddWaiter(waiter);
AddWaiter(&waiter);
// Perform the wait.
return waiter.Wait(lk, timeout);
@ -52,14 +48,14 @@ public:
// Wake up threads in order of priority.
for (auto it = wait_list.begin(); it != wait_list.end();) {
auto& waiter = *it;
if (waiter.need_count > token_count) {
auto* waiter = *it;
if (waiter->need_count > token_count) {
it++;
continue;
}
it = wait_list.erase(it);
token_count -= waiter.need_count;
waiter.cv.notify_one();
token_count -= waiter->need_count;
waiter->cv.notify_one();
}
return true;
@ -70,9 +66,9 @@ public:
if (num_waiters) {
*num_waiters = wait_list.size();
}
for (auto& waiter : wait_list) {
waiter.was_cancled = true;
waiter.cv.notify_one();
for (auto* waiter : wait_list) {
waiter->was_cancled = true;
waiter->cv.notify_one();
}
wait_list.clear();
token_count = set_count < 0 ? init_count : set_count;
@ -80,7 +76,7 @@ public:
}
public:
struct WaitingThread : public ListBaseHook {
struct WaitingThread {
std::condition_variable cv;
u32 priority;
s32 need_count;
@ -132,7 +128,7 @@ public:
}
};
void AddWaiter(WaitingThread& waiter) {
void AddWaiter(WaitingThread* waiter) {
// Insert at the end of the list for FIFO order.
if (is_fifo) {
wait_list.push_back(waiter);
@ -140,16 +136,13 @@ public:
}
// Find the first with priority less then us and insert right before it.
auto it = wait_list.begin();
while (it != wait_list.end() && it->priority > waiter.priority) {
while (it != wait_list.end() && (*it)->priority > waiter->priority) {
it++;
}
wait_list.insert(it, waiter);
}
using WaitingThreads =
boost::intrusive::list<WaitingThread, boost::intrusive::base_hook<ListBaseHook>,
boost::intrusive::constant_time_size<false>>;
WaitingThreads wait_list;
std::list<WaitingThread*> wait_list;
std::string name;
std::atomic<s32> token_count;
std::mutex mutex;

View File

@ -10,6 +10,7 @@
#include "common/thread.h"
#include "core/aerolib/aerolib.h"
#include "core/aerolib/stubs.h"
#include "core/cpu_patches.h"
#include "core/libraries/kernel/memory_management.h"
#include "core/libraries/kernel/thread_management.h"
#include "core/linker.h"
@ -85,6 +86,7 @@ void Linker::Execute() {
// Init primary thread.
Common::SetCurrentThreadName("GAME_MainThread");
Libraries::Kernel::pthreadInitSelfMainThread();
InitializeThreadPatchStack();
InitTlsForThread(true);
// Start shared library modules
@ -104,6 +106,8 @@ void Linker::Execute() {
RunMainEntry(m->GetEntryAddress(), &p, ProgramExitFunc);
}
}
CleanupThreadPatchStack();
}
s32 Linker::LoadModule(const std::filesystem::path& elf_name, bool is_dynamic) {

View File

@ -84,6 +84,8 @@ void Emulator::Run(const std::filesystem::path& file) {
// Applications expect to be run from /app0 so mount the file's parent path as app0.
auto* mnt = Common::Singleton<Core::FileSys::MntPoints>::Instance();
mnt->Mount(file.parent_path(), "/app0");
// Certain games may use /hostapp as well such as CUSA001100
mnt->Mount(file.parent_path(), "/hostapp");
// Loading param.sfo file if exists
std::string id;

View File

@ -130,7 +130,7 @@ void CheatsPatches::setupUI() {
// Call the method to fill the list of cheat files
populateFileListCheats();
QLabel* repositoryLabel = new QLabel("Repository:");
QLabel* repositoryLabel = new QLabel(tr("Repository:"));
repositoryLabel->setAlignment(Qt::AlignLeft);
repositoryLabel->setAlignment(Qt::AlignVCenter);
@ -175,7 +175,8 @@ void CheatsPatches::setupUI() {
int ret = QMessageBox::warning(
this, tr("Delete File"),
QString(tr("Do you want to delete the selected file?\n%1")).arg(selectedFileName),
QString(tr("Do you want to delete the selected file?\\n%1").replace("\\n", "\n"))
.arg(selectedFileName),
QMessageBox::Yes | QMessageBox::No);
if (ret == QMessageBox::Yes) {
@ -1123,7 +1124,7 @@ void CheatsPatches::addPatchesToLayout(const QString& filePath) {
void CheatsPatches::updateNoteTextEdit(const QString& patchName) {
if (m_patchInfos.contains(patchName)) {
const PatchInfo& patchInfo = m_patchInfos[patchName];
QString text = QString(tr("Name:") + " %1\n" + tr("Author:") + " %2\n\n%3")
QString text = QString(tr("Name:") + " %1\n" + tr("Author: ") + "%2\n\n%3")
.arg(patchInfo.name)
.arg(patchInfo.author)
.arg(patchInfo.note);

View File

@ -36,6 +36,7 @@ public:
void downloadCheats(const QString& source, const QString& m_gameSerial,
const QString& m_gameVersion, bool showMessageBox);
void downloadPatches(const QString repository, const bool showMessageBox);
void createFilesJson(const QString& repository);
signals:
void downloadFinished();
@ -58,7 +59,6 @@ private:
void applyCheat(const QString& modName, bool enabled);
void applyPatch(const QString& patchName, bool enabled);
void createFilesJson(const QString& repository);
void uncheckAllCheatCheckBoxes();
void updateNoteTextEdit(const QString& patchName);

View File

@ -7,6 +7,7 @@
#include "about_dialog.h"
#include "cheats_patches.h"
#include "common/io_file.h"
#include "common/string_util.h"
#include "common/version.h"
#include "core/file_format/pkg.h"
#include "core/loader.h"
@ -390,6 +391,8 @@ void MainWindow::CreateConnects() {
nullptr, tr("Download Complete"),
QString(tr("Patches Downloaded Successfully!") + "\n" +
tr("All Patches available for all games have been downloaded.")));
cheatsPatches->createFilesJson("GoldHEN");
cheatsPatches->createFilesJson("shadPS4");
panelDialog->accept();
});
panelDialog->exec();
@ -615,39 +618,48 @@ void MainWindow::InstallDragDropPkg(std::filesystem::path file, int pkgNum, int
pkg = PKG();
pkg.Open(file);
std::string failreason;
const auto extract_path =
std::filesystem::path(Config::getGameInstallDir()) / pkg.GetTitleID();
auto extract_path = std::filesystem::path(Config::getGameInstallDir()) / pkg.GetTitleID();
QString pkgType = QString::fromStdString(pkg.GetPkgFlags());
QDir game_dir(QString::fromStdString(extract_path.string()));
if (game_dir.exists()) {
QMessageBox msgBox;
msgBox.setWindowTitle(tr("PKG Extraction"));
if (pkgType.contains("PATCH")) {
psf.open("", pkg.sfo);
std::string content_id = psf.GetString("CONTENT_ID");
std::string entitlement_label = Common::SplitString(content_id, '-')[2];
auto addon_extract_path = Common::FS::GetUserPath(Common::FS::PathType::AddonsDir) /
pkg.GetTitleID() / entitlement_label;
QDir addon_dir(QString::fromStdString(addon_extract_path.string()));
auto category = psf.GetString("CATEGORY");
if (pkgType.contains("PATCH")) {
QString pkg_app_version = QString::fromStdString(psf.GetString("APP_VER"));
psf.open(extract_path.string() + "/sce_sys/param.sfo", {});
QString game_app_version = QString::fromStdString(psf.GetString("APP_VER"));
double appD = game_app_version.toDouble();
double pkgD = pkg_app_version.toDouble();
if (pkgD == appD) {
msgBox.setText(QString(tr("Patch detected!\nPKG and Game versions match!: "
"%1\nWould you like ") +
tr("to overwrite?"))
.arg(pkg_app_version));
msgBox.setText(QString(tr("Patch detected!") + "\n" +
tr("PKG and Game versions match: ") + pkg_app_version +
"\n" + tr("Would you like to overwrite?")));
msgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
msgBox.setDefaultButton(QMessageBox::No);
} else if (pkgD < appD) {
msgBox.setText(QString(tr("Patch detected!\nPKG Version %1 is older ") +
tr("than installed version!: %2\nWould you like ") +
tr("to overwrite?"))
.arg(pkg_app_version, game_app_version));
msgBox.setText(QString(tr("Patch detected!") + "\n" +
tr("PKG Version %1 is older than installed version: ")
.arg(pkg_app_version) +
game_app_version + "\n" +
tr("Would you like to overwrite?")));
msgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
msgBox.setDefaultButton(QMessageBox::No);
} else {
msgBox.setText(
QString(tr("Patch detected!\nGame is installed: %1\nWould you like ") +
tr("to install Patch: %2?"))
.arg(game_app_version, pkg_app_version));
msgBox.setText(QString(tr("Patch detected!") + "\n" +
tr("Game is installed: ") + game_app_version + "\n" +
tr("Would you like to install Patch: ") +
pkg_app_version + " ?"));
msgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
msgBox.setDefaultButton(QMessageBox::No);
}
@ -657,10 +669,38 @@ void MainWindow::InstallDragDropPkg(std::filesystem::path file, int pkgNum, int
} else {
return;
}
} else if (category == "ac") {
if (!addon_dir.exists()) {
QMessageBox addonMsgBox;
addonMsgBox.setWindowTitle(tr("DLC Installation"));
addonMsgBox.setText(QString(tr("Would you like to install DLC: %1?"))
.arg(QString::fromStdString(entitlement_label)));
addonMsgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
addonMsgBox.setDefaultButton(QMessageBox::No);
int result = addonMsgBox.exec();
if (result == QMessageBox::Yes) {
extract_path = addon_extract_path;
} else {
msgBox.setText(
QString(tr("Game already installed\n%1\nWould you like to overwrite?"))
.arg(QString::fromStdString(extract_path.string())));
return;
}
} else {
msgBox.setText(QString(tr("DLC already installed:") + "\n" +
QString::fromStdString(addon_extract_path.string()) +
"\n\n" + tr("Would you like to overwrite?")));
msgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
msgBox.setDefaultButton(QMessageBox::No);
int result = msgBox.exec();
if (result == QMessageBox::Yes) {
extract_path = addon_extract_path;
} else {
return;
}
}
} else {
msgBox.setText(QString(tr("Game already installed") + "\n" +
QString::fromStdString(extract_path.string()) + "\n" +
tr("Would you like to overwrite?")));
msgBox.setStandardButtons(QMessageBox::Yes | QMessageBox::No);
msgBox.setDefaultButton(QMessageBox::No);
int result = msgBox.exec();
@ -685,6 +725,7 @@ void MainWindow::InstallDragDropPkg(std::filesystem::path file, int pkgNum, int
} else {
int nfiles = pkg.GetNumberOfFiles();
if (nfiles > 0) {
QVector<int> indices;
for (int i = 0; i < nfiles; i++) {
indices.append(i);
@ -725,6 +766,7 @@ void MainWindow::InstallDragDropPkg(std::filesystem::path file, int pkgNum, int
QtConcurrent::map(indices, [&](int index) { pkg.ExtractFiles(index); }));
dialog.exec();
}
}
} else {
QMessageBox::critical(this, tr("PKG ERROR"),
tr("File doesn't appear to be a valid PKG file"));

View File

@ -566,44 +566,54 @@
<translation>PKG-udtrækning</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch opdaget!\nPKG- og spilversioner stemmer overens!: %1\nVil du </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>overskrive?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch opdaget!\nPKG-version %1 er ældre </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>end installeret version!: %2\nVil du </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>overskrive?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Opdatering detekteret!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch opdaget!\nSpillet er installeret: %1\nVil du </translation>
<source>PKG and Game versions match: </source>
<translation>PKG og spilversioner matcher: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>installere patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Vil du overskrive?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Spil allerede installeret\n%1\nVil du overskrive?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG Version %1 er ældre end den installerede version: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Spillet er installeret: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Vil du installere opdateringen: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC Installation</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Vil du installere DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC allerede installeret:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Spillet er allerede installeret</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG-Extraktion</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch erkannt!\nPKG- und Spielversion stimmen überein!: %1\nMöchten Sie </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>überschreiben?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch erkannt!\nPKG-Version %1 ist älter </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>als die installierte Version!: %2\nMöchten Sie </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>überschreiben?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch erkannt!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch erkannt!\nSpiel ist installiert: %1\nMöchten Sie </translation>
<source>PKG and Game versions match: </source>
<translation>PKG- und Spielversionen stimmen überein: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>Patch installieren: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Würden Sie gerne überschreiben?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Spiel bereits installiert\n%1\nMöchten Sie überschreiben?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG-Version %1 ist älter als die installierte Version: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Spiel ist installiert: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Möchten Sie den Patch installieren: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC-Installation</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Würden Sie gerne DLC installieren: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC bereits installiert:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Spiel bereits installiert</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Εξαγωγή PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Ανίχνευση Patch!\nΟι εκδόσεις PKG και παιχνιδιού ταιριάζουν!: %1\nΘέλετε </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>να αντικαταστήσετε;</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Ανίχνευση Patch!\nΗ έκδοση PKG %1 είναι παλαιότερη </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>από την εγκατεστημένη έκδοση!: %2\nΘέλετε </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>να αντικαταστήσετε;</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Αναγνώριση ενημέρωσης!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Ανίχνευση Patch!\nΤο παιχνίδι είναι εγκατεστημένο: %1\nΘέλετε </translation>
<source>PKG and Game versions match: </source>
<translation>Οι εκδόσεις PKG και παιχνιδιού ταιριάζουν: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>να εγκαταστήσετε το Patch: %2;</translation>
<source>Would you like to overwrite?</source>
<translation>Θέλετε να αντικαταστήσετε;</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Το παιχνίδι είναι ήδη εγκατεστημένο\n%1\nΘέλετε να αντικαταστήσετε;</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Η έκδοση PKG %1 είναι παλαιότερη από την εγκατεστημένη έκδοση: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Το παιχνίδι είναι εγκατεστημένο: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Θέλετε να εγκαταστήσετε την ενημέρωση: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Εγκατάσταση DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Θέλετε να εγκαταστήσετε το DLC: %1;</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC ήδη εγκατεστημένο:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Παιχνίδι ήδη εγκατεστημένο</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG Extraction</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch detected!\nPKG and Game versions match!: %1\nWould you like </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>to overwrite?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch detected!\nPKG Version %1 is older </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>than installed version!: %2\nWould you like </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>to overwrite?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch detected!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch detected!\nGame is installed: %1\nWould you like </translation>
<source>PKG and Game versions match: </source>
<translation>PKG and Game versions match: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>to install Patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Would you like to overwrite?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Game already installed\n%1\nWould you like to overwrite?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG Version %1 is older than installed version: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Game is installed: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Would you like to install Patch: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC Installation</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Would you like to install DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC already installed:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Game already installed</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Extracción de PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>¡Parche detectado!\n¡La versión de PKG y del juego coinciden!: %1\n¿Te gustaría </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>¿sobrescribir?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>¡Parche detectado!\nLa versión de PKG %1 es más antigua </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>que la versión instalada!: %2\n¿Te gustaría </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>¿sobrescribir?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>¡Actualización detectada!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>¡Parche detectado!\nJuego está instalado: %1\n¿Te gustaría </translation>
<source>PKG and Game versions match: </source>
<translation>Las versiones de PKG y del juego coinciden: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>¿instalar el parche: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>¿Desea sobrescribir?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Juego ya instalado\n%1\n¿Te gustaría sobrescribirlo?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>La versión de PKG %1 es más antigua que la versión instalada: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>El juego está instalado: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>¿Desea instalar la actualización: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Instalación de DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>¿Desea instalar el DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC ya instalado:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Juego ya instalado</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG:n purku</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Korjaus havaittu!\nPKG:n ja pelin versiot vastaavat!: %1\nHaluatko </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>korvata?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Korjaus havaittu!\nPKG Version %1 on vanhempi </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>kuin asennettu versio!: %2\nHaluatko </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>korvata?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Päivitys havaittu!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Korjaus havaittu!\nPeli on asennettu: %1\nHaluatko </translation>
<source>PKG and Game versions match: </source>
<translation>PKG- ja peliversiot vastaavat: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>asentaa korjaus: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Haluatko korvata?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Peli on jo asennettu\n%1\nHaluatko korvata sen?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG-versio %1 on vanhempi kuin asennettu versio: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Peli on asennettu: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Haluatko asentaa päivityksen: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC-asennus</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Haluatko asentaa DLC:n: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC on jo asennettu:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Peli on jo asennettu</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Extraction du PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch détecté !\nLa version du PKG et du jeu correspondent : %1\nSouhaitez-vous </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>écraser ?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch détecté !\nVersion PKG %1 est plus ancienne </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>que la version installée ! : %2\nSouhaitez-vous </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>écraser ?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch détecté !</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch détecté !\nJeu est installé : %1\nSouhaitez-vous </translation>
<source>PKG and Game versions match: </source>
<translation>Les versions PKG et jeu correspondent : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>installer le patch : %2 ?</translation>
<source>Would you like to overwrite?</source>
<translation>Souhaitez-vous remplacer ?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Jeu déjà installé\n%1\nSouhaitez-vous écraser ?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>La version PKG %1 est plus ancienne que la version installée : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Jeu installé : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Souhaitez-vous installer le patch : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Installation du DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Souhaitez-vous installer le DLC : %1 ?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC déjà installé :</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Jeu déjà installé</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG kicsomagolás</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Javítás észlelve!\nA PKG és a játék verziók egyeznek: %1\nSzeretnéd </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>felülírni?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Javítás észlelve!\nA PKG verzió %1 régebbi </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>mint a telepített verzió: %2\nSzeretnéd </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>felülírni?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Frissítés észlelve!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Javítás észlelve!\nA játék telepítve van: %1\nSzeretnéd </translation>
<source>PKG and Game versions match: </source>
<translation>A PKG és a játék verziói egyeznek: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>a javítást telepíteni: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Szeretné felülírni?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>A játék már telepítve van\n%1\nSzeretnéd felülírni?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>A %1-es PKG verzió régebbi, mint a telepített verzió: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>A játék telepítve van: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Szeretné telepíteni a frissítést: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC Telepítés</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Szeretné telepíteni a DLC-t: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC már telepítve:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>A játék már telepítve van</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Ekstraksi PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch terdeteksi!\nVersi PKG dan Game cocok!: %1\nApakah Anda ingin </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>menimpa?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch terdeteksi!\nVersi PKG %1 lebih lama </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>daripada versi yang terpasang!: %2\nApakah Anda ingin </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>menimpa?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch terdeteksi!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch terdeteksi!\nGame terpasang: %1\nApakah Anda ingin </translation>
<source>PKG and Game versions match: </source>
<translation>Versi PKG dan Game cocok: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>memasang Patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Apakah Anda ingin menimpa?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Game sudah terpasang\n%1\nApakah Anda ingin menimpa?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Versi PKG %1 lebih lama dari versi yang terpasang: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Game telah terpasang: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Apakah Anda ingin menginstal patch: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Instalasi DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Apakah Anda ingin menginstal DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC sudah terpasang:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Game sudah terpasang</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -181,7 +181,7 @@
<message>
<location filename="../main_window_ui.h" line="318"/>
<source>Install application from a .pkg file</source>
<translation>Installa applicazione da un file .pkg file</translation>
<translation>Installa applicazione da un file .pkg</translation>
</message>
<message>
<location filename="../main_window_ui.h" line="320"/>
@ -236,7 +236,7 @@
<message>
<location filename="../main_window_ui.h" line="338"/>
<source>List View</source>
<translation>Visualizzazione lista</translation>
<translation>Visualizzazione Lista</translation>
</message>
<message>
<location filename="../main_window_ui.h" line="340"/>
@ -341,7 +341,7 @@
<message>
<location filename="../main_window_ui.h" line="364"/>
<source>toolBar</source>
<translation>barra strumenti</translation>
<translation>Barra strumenti</translation>
</message>
</context>
<context>
@ -356,8 +356,8 @@
<name>TrophyViewer</name>
<message>
<location filename="../trophy_viewer.cpp" line="8"/>
<source>Visualizzatore Trofei</source>
<translation>Trophy Viewer</translation>
<source>Trophy Viewer</source>
<translation>Visualizzatore Trofei</translation>
</message>
</context>
<context>
@ -503,42 +503,42 @@
<message>
<location filename="../main_window.cpp" line="168"/>
<source> * Unsupported Vulkan Version</source>
<translation> * Versi Vulkan Tidak Didukung</translation>
<translation> * Versione Vulkan non supportata</translation>
</message>
<message>
<location filename="../main_window.cpp" line="326"/>
<source>Download Cheats For All Installed Games</source>
<translation>Unduh Cheat Untuk Semua Game yang Terinstal</translation>
<translation>Scarica Trucchi per tutti i giochi installati</translation>
</message>
<message>
<location filename="../main_window.cpp" line="328"/>
<source>Download Patches For All Games</source>
<translation>Unduh Patch Untuk Semua Game</translation>
<translation>Scarica Patch per tutti i giochi</translation>
</message>
<message>
<location filename="../main_window.cpp" line="363"/>
<source>Download Complete</source>
<translation>Unduhan Selesai</translation>
<translation>Scaricamento completato</translation>
</message>
<message>
<location filename="../main_window.cpp" line="364"/>
<source>You have downloaded cheats for all the games you have installed.</source>
<translation>Anda telah mengunduh cheat untuk semua game yang telah Anda instal.</translation>
<translation>Hai scaricato trucchi per tutti i giochi installati.</translation>
</message>
<message>
<location filename="../main_window.cpp" line="391"/>
<source>Patches Downloaded Successfully!</source>
<translation>Patch Berhasil Diunduh!</translation>
<translation>Patch scaricate con successo!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="392"/>
<source>All Patches available for all games have been downloaded.</source>
<translation>Semua patch yang tersedia untuk semua game telah diunduh.</translation>
<translation>Tutte le patch disponibili per tutti i giochi sono state scaricate.</translation>
</message>
<message>
<location filename="../main_window.cpp" line="549"/>
<source>Games: </source>
<translation>Game: </translation>
<translation>Giochi: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="575"/>
@ -553,87 +553,97 @@
<message>
<location filename="../main_window.cpp" line="600"/>
<source>Game Boot</source>
<translation>Boot Game</translation>
<translation>Avvia Gioco</translation>
</message>
<message>
<location filename="../main_window.cpp" line="600"/>
<source>Only one file can be selected!</source>
<translation>Hanya satu file yang dapat dipilih!</translation>
<translation>Si può selezionare solo un file!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="623"/>
<source>PKG Extraction</source>
<translation>Ekstraksi PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch terdeteksi!\nVersi PKG dan Game cocok!: %1\nApakah Anda ingin </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>menimpa?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch terdeteksi!\nVersi PKG %1 lebih lama </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>daripada versi yang terinstal!: %2\nApakah Anda ingin </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>menimpa?</translation>
<translation>Estrazione file PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch terdeteksi!\nGame terinstal: %1\nApakah Anda ingin </translation>
<source>Patch detected!</source>
<translation>Patch rilevata!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>PKG and Game versions match: </source>
<translation>Le versioni di PKG e del gioco corrispondono: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>menginstal Patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Vuoi sovrascrivere?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Game sudah terinstal\n%1\nApakah Anda ingin menimpa?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>La versione PKG %1 è più vecchia rispetto alla versione installata: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Gioco installato: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Vuoi installare la patch: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Installazione DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Vuoi installare il DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC già installato:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Gioco già installato</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>PKG is a patch, please install the game first!</source>
<translation>PKG adalah patch, silakan instal game terlebih dahulu!</translation>
<translation>Questo file PKG contiene una patch. Per favore, installa prima il gioco!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="681"/>
<source>PKG ERROR</source>
<translation>ERROR PKG</translation>
<translation>ERRORE PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="693"/>
<source>Extracting PKG %1/%2</source>
<translation>Estrazione PKG %1/%2</translation>
<translation>Estrazione file PKG %1/%2</translation>
</message>
<message>
<location filename="../main_window.cpp" line="703"/>
<source>Extraction Finished</source>
<translation>Ekstraksi Selesai</translation>
<translation>Estrazione Completata</translation>
</message>
<message>
<location filename="../main_window.cpp" line="704"/>
<source>Game successfully installed at %1</source>
<translation>Game berhasil diinstal di %1</translation>
<translation>Gioco installato correttamente in %1</translation>
</message>
<message>
<location filename="../main_window.cpp" line="725"/>
<source>File doesn't appear to be a valid PKG file</source>
<translation>File tidak tampak sebagai file PKG yang valid</translation>
<translation>Il file sembra non essere un file PKG valido</translation>
</message>
</context>
<context>
@ -641,87 +651,87 @@
<message>
<location filename="../cheats_patches.cpp" line="44"/>
<source>Cheats / Patches</source>
<translation>Cheat / Patch</translation>
<translation>Trucchi / Patch</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="50"/>
<source>defaultTextEdit_MSG</source>
<translation>I cheats/patches sono sperimentali.\nUtilizzali con cautela.\n\nScarica i cheats singolarmente selezionando il repository e cliccando sul pulsante di download.\nNella scheda Patches, puoi scaricare tutti i patch in una volta sola, scegliere quali vuoi utilizzare e salvare la tua selezione.\n\nPoiché non sviluppiamo i Cheats/Patches,\nper favore segnala i problemi all'autore del cheat.\n\nHai creato un nuovo cheat? Visita:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
<translation>I trucchi e le patch sono sperimentali.\nUtilizzali con cautela.\n\nScarica i trucchi singolarmente selezionando l'archivio e cliccando sul pulsante di download.\nNella scheda Patch, puoi scaricare tutte le patch in una volta sola, scegliere quali vuoi utilizzare e salvare la tua selezione.\n\nPoiché non sviluppiamo i trucchi e le patch,\nper favore segnala i problemi all'autore dei trucchi.\n\nHai creato un nuovo trucco? Visita:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="69"/>
<source>No Image Available</source>
<translation>Tidak Ada Gambar</translation>
<translation>Nessuna immagine disponibile</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="79"/>
<source>Serial: </source>
<translation>Serial: </translation>
<translation>Seriale: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="83"/>
<source>Version: </source>
<translation>Versi: </translation>
<translation>Versione: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="87"/>
<source>Size: </source>
<translation>Ukuran: </translation>
<translation>Dimensione: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="126"/>
<source>Select Cheat File:</source>
<translation>Pilih File Cheat:</translation>
<translation>Seleziona File Trucchi:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="133"/>
<source>Repository:</source>
<translation>Repositori:</translation>
<translation>Archivio:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="149"/>
<source>Download Cheats</source>
<translation>Unduh Cheat</translation>
<translation>Scarica trucchi</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="155"/>
<source>Delete File</source>
<translation>Hapus File</translation>
<translation>Cancella File</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="169"/>
<source>No files selected.</source>
<translation>Tidak ada file yang dipilih.</translation>
<translation>Nessun file selezionato.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="170"/>
<source>You can delete the cheats you don't want after downloading them.</source>
<translation>Anda dapat menghapus cheat yang tidak diinginkan setelah mengunduhnya.</translation>
<translation>Puoi cancellare i trucchi che non vuoi utilizzare dopo averli scaricati.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="178"/>
<source>Do you want to delete the selected file?\n%1</source>
<translation>Apakah Anda ingin menghapus file yang dipilih?\n%1</translation>
<translation>Vuoi cancellare il file selezionato?\n%1</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="213"/>
<source>Select Patch File:</source>
<translation>Pilih File Patch:</translation>
<translation>Seleziona File Patch:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="241"/>
<source>Download Patches</source>
<translation>Unduh Patch</translation>
<translation>Scarica Patch</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="248"/>
<source>Save</source>
<translation>Simpan</translation>
<translation>Salva</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="256"/>
<source>Cheats</source>
<translation>Cheat</translation>
<translation>Trucchi</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="257"/>
@ -731,162 +741,162 @@
<message>
<location filename="../cheats_patches.cpp" line="278"/>
<source>Error</source>
<translation>Kesalahan</translation>
<translation>Errore</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="278"/>
<source>No patch selected.</source>
<translation>Tidak ada patch yang dipilih.</translation>
<translation>Nessuna patch selezionata.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="292"/>
<source>Unable to open files.json for reading.</source>
<translation>Gagal membuka files.json untuk dibaca.</translation>
<translation>Impossibile aprire il file .json per la lettura.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="316"/>
<source>No patch file found for the current serial.</source>
<translation>Tidak ada file patch ditemukan untuk serial saat ini.</translation>
<translation>Nessun file patch trovato per il seriale selezionato.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="323"/>
<source>Unable to open the file for reading.</source>
<translation>Gagal membuka file untuk dibaca.</translation>
<translation>Impossibile aprire il file per la lettura.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="422"/>
<source>Unable to open the file for writing.</source>
<translation>Gagal membuka file untuk ditulis.</translation>
<translation>Impossibile aprire il file per la scrittura.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="432"/>
<source>Failed to parse XML: </source>
<translation>Gagal mengurai XML: </translation>
<translation>Analisi XML fallita: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="434"/>
<source>Success</source>
<translation>Berhasil</translation>
<translation>Successo</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="434"/>
<source>Options saved successfully.</source>
<translation>Opsi berhasil disimpan.</translation>
<translation>Opzioni salvate con successo.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="470"/>
<source>Invalid Source</source>
<translation>Sumber Tidak Valid</translation>
<translation>Fonte non valida</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="471"/>
<source>The selected source is invalid.</source>
<translation>Sumber yang dipilih tidak valid.</translation>
<translation>La fonte selezionata non è valida.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="519"/>
<source>File Exists</source>
<translation>File Ada</translation>
<translation>Il file è presente</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="520"/>
<source>File already exists. Do you want to replace it?</source>
<translation>File sudah ada. Apakah Anda ingin menggantinya?</translation>
<translation>Il file è già presente. Vuoi sostituirlo?</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="539"/>
<source>Failed to save file:</source>
<translation>Gagal menyimpan file:</translation>
<translation>Salvataggio file fallito:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="545"/>
<source>Failed to download file:</source>
<translation>Gagal mengunduh file:</translation>
<translation>Scaricamento file fallito:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="556"/>
<source>Cheats Not Found</source>
<translation>Cheat Tidak Ditemukan</translation>
<translation>Trucchi non trovati</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="556"/>
<source>CheatsNotFound_MSG</source>
<translation>Cheat tidak ditemukan untuk game ini dalam versi repositori yang dipilih, coba repositori lain atau versi game yang berbeda.</translation>
<translation>Non sono stati trovati trucchi per questa versione del gioco nell'archivio selezionato, prova un altro archivio o una versione diversa del gioco.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="593"/>
<source>Cheats Downloaded Successfully</source>
<translation>Cheat Berhasil Diunduh</translation>
<translation>Trucchi scaricati con successo!</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="594"/>
<source>CheatsDownloadedSuccessfully_MSG</source>
<translation>Anda telah berhasil mengunduh cheat untuk versi game ini dari repositori yang dipilih. Anda dapat mencoba mengunduh dari repositori lain, jika tersedia, Anda juga dapat menggunakannya dengan memilih file dari daftar.</translation>
<translation>Hai scaricato con successo i trucchi per questa versione del gioco dall'archivio selezionato. Puoi provare a scaricare da un altro archivio, se disponibile, puoi anche utilizzarlo selezionando il file dall'elenco.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="747"/>
<source>Failed to save:</source>
<translation>Gagal menyimpan:</translation>
<translation>Salvataggio fallito:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="754"/>
<source>Failed to download:</source>
<translation>Gagal mengunduh:</translation>
<translation>Impossibile scaricare:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="762"/>
<source>Download Complete</source>
<translation>Unduhan Selesai</translation>
<translation>Scaricamento completo</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="763"/>
<source>DownloadComplete_MSG</source>
<translation>Patch Berhasil Diunduh! Semua patch yang tersedia untuk semua game telah diunduh, tidak perlu mengunduhnya secara individu untuk setiap game seperti yang terjadi pada Cheat.</translation>
<translation>Patch scaricata con successo! Vengono scaricate tutte le patch disponibili per tutti i giochi, non è necessario scaricarle singolarmente per ogni gioco come nel caso dei trucchi.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="773"/>
<source>Failed to parse JSON data from HTML.</source>
<translation>Gagal mengurai data JSON dari HTML.</translation>
<translation>Impossibile analizzare i dati JSON dall'HTML.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="778"/>
<source>Failed to retrieve HTML page.</source>
<translation>Gagal mengambil halaman HTML.</translation>
<translation>Impossibile recuperare la pagina HTML.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="801"/>
<source>Failed to open file:</source>
<translation>Gagal membuka file:</translation>
<translation>Impossibile aprire file:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="819"/>
<source>XML ERROR:</source>
<translation>KESALAHAN XML:</translation>
<translation>ERRORE XML:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="826"/>
<source>Failed to open files.json for writing</source>
<translation>Gagal membuka files.json untuk menulis</translation>
<translation>Impossibile aprire i file .json per la scrittura</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="925"/>
<source>Author: </source>
<translation>Penulis: </translation>
<translation>Autore: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="997"/>
<source>Directory does not exist:</source>
<translation>Direktori tidak ada:</translation>
<translation>La cartella non esiste:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="1006"/>
<source>Failed to open files.json for reading.</source>
<translation>Gagal membuka files.json untuk dibaca.</translation>
<translation>Impossibile aprire i file .json per la lettura.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="1006"/>
<source>Name:</source>
<translation>Nama:</translation>
<translation>Nome:</translation>
</message>
</context>
</TS>

View File

@ -566,44 +566,54 @@
<translation>PKG抽出</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>\nPKGとゲームバージョンが一致しています: %1\n上書きしますか</translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>\nPKGバージョン %1 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>: %2\n上書きしますか</translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation></translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>\nゲームがインストールされています: %1\nインストールしますか</translation>
<source>PKG and Game versions match: </source>
<translation>PKGとゲームのバージョンが一致しています: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>: %2?</translation>
<source>Would you like to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>\n%1\n上書きしますか</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKGバージョン %1 : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLCのインストール</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>DLCをインストールしてもよろしいですか: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLCはすでにインストールされています:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG Extraction</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch detected!\nPKG and Game versions match!: %1\nWould you like </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>to overwrite?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch detected!\nPKG Version %1 is older </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>than installed version!: %2\nWould you like </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>to overwrite?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch detected!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch detected!\nGame is installed: %1\nWould you like </translation>
<source>PKG and Game versions match: </source>
<translation>PKG and Game versions match: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>to install Patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Would you like to overwrite?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Game already installed\n%1\nWould you like to overwrite?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG Version %1 is older than installed version: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Game is installed: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Would you like to install Patch: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC Installation</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Would you like to install DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC already installed:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Game already installed</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG ištraukimas</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Pataisa aptikta!\nPKG ir žaidimo versijos atitinka!: %1\nAr norėtumėte </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>perrašyti?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Pataisa aptikta!\nPKG versija %1 yra senesnė </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>nei įdiegta versija!: %2\nAr norėtumėte </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>perrašyti?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Rasta atnaujinimą!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Pataisa aptikta!\nŽaidimas įdiegtas: %1\nAr norėtumėte </translation>
<source>PKG and Game versions match: </source>
<translation>PKG ir žaidimo versijos sutampa: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>įdiegti pataisą: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Ar norite perrašyti?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Žaidimas jau įdiegtas\n%1\nAr norėtumėte perrašyti?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG versija %1 yra senesnė nei įdiegta versija: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Žaidimas įdiegtas: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Ar norite įdiegti atnaujinimą: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC diegimas</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Ar norite įdiegti DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC jau įdiegtas:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Žaidimas jau įdiegtas</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG-ekstraksjon</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Oppdatering oppdaget!\nPKG og spillversjoner stemmer!: %1\nØnsker du å </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>overskrive?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Oppdatering oppdaget!\nPKG-versjon %1 er eldre </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>enn installert versjon!: %2\nØnsker du å </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>overskrive?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Oppdatering oppdaget!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Oppdatering oppdaget!\nSpillet er installert: %1\nØnsker du å </translation>
<source>PKG and Game versions match: </source>
<translation>PKG- og spillversjoner stemmer overens: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>installere oppdateringen: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Ønsker du å overskrive?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Spill allerede installert\n%1\nØnsker du å overskrive?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG-versjon %1 er eldre enn installert versjon: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Spillet er installert: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Ønsker du å installere oppdateringen: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC-installasjon</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Ønsker du å installere DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC allerede installert:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Spillet er allerede installert</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG-extractie</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch gedetecteerd!\nPKG en spelversies komen overeen!: %1\nWil je </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>overschrijven?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch gedetecteerd!\nPKG-versie %1 is ouder </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>dan de geïnstalleerde versie!: %2\nWil je </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>overschrijven?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch gedetecteerd!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch gedetecteerd!\nSpel is geïnstalleerd: %1\nWil je </translation>
<source>PKG and Game versions match: </source>
<translation>PKG- en gameversies komen overeen: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>de patch installeren: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Wilt u overschrijven?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Spel al geïnstalleerd\n%1\nWil je het overschrijven?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG-versie %1 is ouder dan de geïnstalleerde versie: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Game is geïnstalleerd: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Wilt u de patch installeren: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC-installatie</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Wilt u DLC installeren: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC al geïnstalleerd:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Game al geïnstalleerd</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Ekstrakcja PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Wykryto poprawkę!\nWersje PKG i gry pasują do siebie!: %1\nCzy chcesz </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>nadpisać?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Wykryto poprawkę!\nWersja PKG %1 jest starsza </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>niż zainstalowana wersja!: %2\nCzy chcesz </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>nadpisać?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Wykryto łatkę!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Wykryto poprawkę!\nGra jest zainstalowana: %1\nCzy chcesz </translation>
<source>PKG and Game versions match: </source>
<translation>Wersje PKG i gry zgodne: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>zainstalować poprawkę: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Czy chcesz nadpisać?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Gra już zainstalowana\n%1\nCzy chcesz nadpisać?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Wersja PKG %1 jest starsza niż zainstalowana wersja: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Gra jest zainstalowana: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Czy chcesz zainstalować łatkę: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Instalacja DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Czy chcesz zainstalować DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC już zainstalowane:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Gra już zainstalowana</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -37,7 +37,7 @@
<message>
<location filename="../game_info.cpp" line="26"/>
<source>Loading game list, please wait :3</source>
<translation>Carregando lista de jogos, por favor aguarde :3</translation>
<translation>Carregando a lista de jogos, por favor aguarde :3</translation>
</message>
<message>
<location filename="../game_info.cpp" line="26"/>
@ -256,7 +256,7 @@
<message>
<location filename="../main_window_ui.h" line="343"/>
<source>Download Cheats/Patches</source>
<translation>Baixar Trapaças / Patches</translation>
<translation>Baixar Cheats/Patches</translation>
</message>
<message>
<location filename="../main_window_ui.h" line="345"/>
@ -291,7 +291,7 @@
<message>
<location filename="../main_window_ui.h" line="354"/>
<source>Game List Mode</source>
<translation>Modo de Lista de Jogos</translation>
<translation>Modo da Lista de Jogos</translation>
</message>
<message>
<location filename="../main_window_ui.h" line="355"/>
@ -425,7 +425,7 @@
<message>
<location filename="../settings_dialog.ui" line="235"/>
<source>Log Filter</source>
<translation>Filtro</translation>
<translation>Filtro do Registro</translation>
</message>
<message>
<location filename="../settings_dialog.ui" line="272"/>
@ -508,12 +508,12 @@
<message>
<location filename="../main_window.cpp" line="326"/>
<source>Download Cheats For All Installed Games</source>
<translation>Baixar Trapaças para todos os jogos instalados</translation>
<translation>Baixar Cheats para Todos os Jogos Instalados</translation>
</message>
<message>
<location filename="../main_window.cpp" line="328"/>
<source>Download Patches For All Games</source>
<translation>Baixar Patches para todos os jogos</translation>
<translation>Baixar Patches para Todos os Jogos</translation>
</message>
<message>
<location filename="../main_window.cpp" line="363"/>
@ -523,7 +523,7 @@
<message>
<location filename="../main_window.cpp" line="364"/>
<source>You have downloaded cheats for all the games you have installed.</source>
<translation>Você baixou trapaças para todos os jogos que instalou.</translation>
<translation>Você baixou cheats para todos os jogos que instalou.</translation>
</message>
<message>
<location filename="../main_window.cpp" line="391"/>
@ -566,54 +566,64 @@
<translation>Extração de PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch detectado!\nVersões PKG e do Jogo correspondem!: %1\nGostaria de </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>substituir?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch detectado!\nVersão PKG %1 é mais antiga </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>do que a versão instalada!: %2\nGostaria de </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>substituir?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Atualização detectada!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch detectado!\nJogo está instalado: %1\nGostaria de </translation>
<source>PKG and Game versions match: </source>
<translation>As versões do PKG e do Jogo são igual: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>instalar o Patch: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Gostaria de substituir?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Jogo instalado\n%1\nGostaria de substituir?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Versão do PKG %1 é mais antiga do que a versão instalada: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Jogo instalado: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Você gostaria de instalar a atualização: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Instalação de DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Você gostaria de instalar o DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC instalada:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>O jogo está instalado:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>PKG is a patch, please install the game first!</source>
<translation>PKG é um patch, por favor, instale o jogo primeiro!</translation>
<translation>O PKG é um patch, por favor, instale o jogo primeiro!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="681"/>
<source>PKG ERROR</source>
<translation>ERRO PKG</translation>
<translation>ERRO de PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="693"/>
@ -641,12 +651,12 @@
<message>
<location filename="../cheats_patches.cpp" line="44"/>
<source>Cheats / Patches</source>
<translation>Trapaças / Patches</translation>
<translation>Cheats / Patches</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="50"/>
<source>defaultTextEdit_MSG</source>
<translation>Trapaças/Patches são experimentais.\nUse com cautela.\n\nBaixe as trapaças individualmente selecionando o repositório e clicando no botão de download.\nNa aba Patches, você pode baixar todos os Patches de uma vez, escolher qual deseja usar e salvar a opção.\n\nComo não desenvolvemos as Trapaças/Patches,\npor favor, reporte problemas relacionados ao autor da trapaça.\n\nCriou uma nova trapaça? Visite:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
<translation>Cheats/Patches são experimentais.\nUse com cautela.\n\nBaixe os cheats individualmente selecionando o repositório e clicando no botão de download.\nNa aba Patches, você pode baixar todos os Patches de uma vez, escolha qual deseja usar e salve a opção.\n\nComo não desenvolvemos os Cheats/Patches,\npor favor, reporte problemas relacionados ao autor do cheat.\n\nCriou um novo cheat? Visite:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="69"/>
@ -656,7 +666,7 @@
<message>
<location filename="../cheats_patches.cpp" line="79"/>
<source>Serial: </source>
<translation>Série: </translation>
<translation>Serial: </translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="83"/>
@ -671,7 +681,7 @@
<message>
<location filename="../cheats_patches.cpp" line="126"/>
<source>Select Cheat File:</source>
<translation>Selecione o Arquivo de Trapaça:</translation>
<translation>Selecione o Arquivo de Cheat:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="133"/>
@ -681,7 +691,7 @@
<message>
<location filename="../cheats_patches.cpp" line="149"/>
<source>Download Cheats</source>
<translation>Baixar Trapaças</translation>
<translation>Baixar Cheats</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="155"/>
@ -696,7 +706,7 @@
<message>
<location filename="../cheats_patches.cpp" line="170"/>
<source>You can delete the cheats you don't want after downloading them.</source>
<translation>Você pode excluir as trapaças que não deseja após baixá-las.</translation>
<translation>Você pode excluir os cheats que não deseja após baixá-las.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="178"/>
@ -721,7 +731,7 @@
<message>
<location filename="../cheats_patches.cpp" line="256"/>
<source>Cheats</source>
<translation>Trapaças</translation>
<translation>Cheats</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="257"/>
@ -746,7 +756,7 @@
<message>
<location filename="../cheats_patches.cpp" line="316"/>
<source>No patch file found for the current serial.</source>
<translation>Nenhum arquivo de patch encontrado para a série atual.</translation>
<translation>Nenhum arquivo de patch encontrado para o serial atual.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="323"/>
@ -806,22 +816,22 @@
<message>
<location filename="../cheats_patches.cpp" line="556"/>
<source>Cheats Not Found</source>
<translation>Trapaças Não Encontradas</translation>
<translation>Cheats Não Encontrados</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="556"/>
<source>CheatsNotFound_MSG</source>
<translation>Nenhuma trapaça encontrada para este jogo nesta versão do repositório selecionado, tente outro repositório ou uma versão diferente do jogo.</translation>
<translation>Nenhum cheat encontrado para este jogo nesta versão do repositório selecionado, tente outro repositório ou uma versão diferente do jogo.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="593"/>
<source>Cheats Downloaded Successfully</source>
<translation>Trapaças Baixadas com Sucesso</translation>
<translation>Cheats Baixados com Sucesso</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="594"/>
<source>CheatsDownloadedSuccessfully_MSG</source>
<translation>Você baixou as trapaças com sucesso. Para esta versão do jogo a partir do repositório selecionado.Você pode tentar baixar de outro repositório, se estiver disponível, também será possível usá-lo selecionando o arquivo da lista.</translation>
<translation>Você baixou os cheats com sucesso. Para esta versão do jogo a partir do repositório selecionado. Você pode tentar baixar de outro repositório, se estiver disponível, também será possível usá-lo selecionando o arquivo da lista.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="747"/>
@ -841,7 +851,7 @@
<message>
<location filename="../cheats_patches.cpp" line="763"/>
<source>DownloadComplete_MSG</source>
<translation>Patches Baixados com Sucesso! Todos os patches disponíveis para todos os jogos foram baixados, não é necessário baixá-los individualmente para cada jogo como acontece com as Trapaças.</translation>
<translation>Patches Baixados com Sucesso! Todos os patches disponíveis para todos os jogos foram baixados, não é necessário baixá-los individualmente para cada jogo como acontece com os Cheats.</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="773"/>
@ -861,7 +871,7 @@
<message>
<location filename="../cheats_patches.cpp" line="819"/>
<source>XML ERROR:</source>
<translation>ERRO XML:</translation>
<translation>ERRO de XML:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="826"/>
@ -876,7 +886,7 @@
<message>
<location filename="../cheats_patches.cpp" line="997"/>
<source>Directory does not exist:</source>
<translation>Diretório não existe:</translation>
<translation>O Diretório não existe:</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="1006"/>

View File

@ -566,44 +566,54 @@
<translation>Extracție PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Patch detectat!\nVersiunile PKG și Joc se potrivesc!: %1\nAi dori </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation> suprascrii?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Patch detectat!\nVersiunea PKG %1 este mai veche </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>decât versiunea instalată!: %2\nAi dori </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation> suprascrii?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Patch detectat!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Patch detectat!\nJocul este instalat: %1\nAi dori </translation>
<source>PKG and Game versions match: </source>
<translation>Versiunile PKG și ale jocului sunt compatibile: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation> instalezi Patch-ul: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Doriți suprascrieți?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Jocul este deja instalat\n%1\nAi dori suprascrii?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Versiunea PKG %1 este mai veche decât versiunea instalată: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Jocul este instalat: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Doriți instalați patch-ul: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Instalare DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Doriți instalați DLC-ul: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC deja instalat:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Jocul deja instalat</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -93,7 +93,7 @@
<message>
<location filename="../gui_context_menus.h" line="48"/>
<source>Cheats / Patches</source>
<translation>Читы / Патчи</translation>
<translation>Читы и патчи</translation>
</message>
<message>
<location filename="../gui_context_menus.h" line="49"/>
@ -256,7 +256,7 @@
<message>
<location filename="../main_window_ui.h" line="343"/>
<source>Download Cheats/Patches</source>
<translation>Скачать Читы / Патчи</translation>
<translation>Скачать читы или патчи</translation>
</message>
<message>
<location filename="../main_window_ui.h" line="345"/>
@ -566,44 +566,54 @@
<translation>Извлечение PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Обнаружен патч!\nВерсии PKG и игры совпадают!: %1\nХотите </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>перезаписать?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Обнаружен патч!\nВерсия PKG %1 устарела </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>по сравнению с установленной версией!: %2\nХотите </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>перезаписать?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Обнаружен патч!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Обнаружен патч!\гра установлена: %1\nХотите </translation>
<source>PKG and Game versions match: </source>
<translation>Версии PKG и игры совпадают: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>установить патч: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Хотите перезаписать?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Игра уже установлена\n%1\nХотите перезаписать?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Версия PKG %1 старее установленной версии: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Игра установлена: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Хотите установить патч: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Установка DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Вы хотите установить DLC: %1??</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC уже установлен:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Игра уже установлена</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
@ -628,7 +638,7 @@
<message>
<location filename="../main_window.cpp" line="704"/>
<source>Game successfully installed at %1</source>
<translation>Игра успешно установлена по адресу %1</translation>
<translation>Игра успешно установлена в %1</translation>
</message>
<message>
<location filename="../main_window.cpp" line="725"/>
@ -641,12 +651,12 @@
<message>
<location filename="../cheats_patches.cpp" line="44"/>
<source>Cheats / Patches</source>
<translation>Читы / Патчи</translation>
<translation>Читы и патчи</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="50"/>
<source>defaultTextEdit_MSG</source>
<translation>Cheats/Patches sunt experimentale.\nUtilizați cu prudență.\n\nDescărcați cheats individual prin selectarea depozitului și făcând clic pe butonul de descărcare.\nÎn fila Patches, puteți descărca toate patch-urile deodată, alege pe cele pe care doriți le utilizați și salvați selecția.\n\nDeoarece nu dezvoltăm Cheats/Patches,\nte rugăm raportezi problemele autorului cheat-ului.\n\nAi creat un nou cheat? Vizitează:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
<translation>Читы и патчи экспериментальны.\спользуйте с осторожностью.\n\nСкачивайте читы, выбрав репозиторий и нажав на кнопку загрузки.\nВо вкладке "Патчи" вы можете скачать все патчи сразу, выбирать какие вы хотите использовать, и сохранять выбор.\n\оскольку мы не разрабатываем читы/патчи,\nпожалуйста сообщайте о проблемах автору чита/патча.\n\nСоздали новый чит? Посетите:\nhttps://github.com/shadps4-emu/ps4_cheats</translation>
</message>
<message>
<location filename="../cheats_patches.cpp" line="69"/>

View File

@ -566,44 +566,54 @@
<translation>PKG Çıkartma</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Yama tespit edildi!\nPKG ve Oyun sürümleri uyuyor!: %1\nÜzerine yazmak ister misiniz?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>üzerine yazmak?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Yama tespit edildi!\nPKG Sürümü %1 daha eski </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>yüklü sürümden!: %2\nÜzerine yazmak ister misiniz?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>üzerine yazmak?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Yamanın tespit edildi!</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Yama tespit edildi!\nOyun yüklü: %1\nÜzerine yazmak ister misiniz?</translation>
<source>PKG and Game versions match: </source>
<translation>PKG ve oyun sürümleri uyumlu: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>Yamayı kurmak ister misiniz: %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Üzerine yazmak ister misiniz?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Oyun zaten yüklü\n%1\nÜzerine yazmak ister misiniz?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG Sürümü %1, kurulu sürümden daha eski: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Oyun yüklendi: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Yamanın yüklenmesini ister misiniz: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC Yükleme</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>DLC'yi yüklemek ister misiniz: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC zaten yüklü:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Oyun zaten yüklü</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>Giải nén PKG</translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>Đã phát hiện bản !\nPhiên bản PKG trò chơi khớp!: %1\nBạn muốn </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation>ghi đè không?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>Đã phát hiện bản !\nPhiên bản PKG %1 hơn </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>so với phiên bản đã cài đt!: %2\nBạn muốn </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation>ghi đè không?</translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation>Đã phát hiện bản !</translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>Đã phát hiện bản !\nTrò chơi đã đưc cài đt: %1\nBạn muốn </translation>
<source>PKG and Game versions match: </source>
<translation>Các phiên bản PKG trò chơi khớp nhau: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>cài đt bản : %2?</translation>
<source>Would you like to overwrite?</source>
<translation>Bạn muốn ghi đè không?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>Trò chơi đã đưc cài đt\n%1\nBạn muốn ghi đè không?</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>Phiên bản PKG %1 hơn phiên bản đã cài đt: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>Trò chơi đã đưc cài đt: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>Bạn muốn cài đt bản : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>Cài đt DLC</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation>Bạn muốn cài đt DLC: %1?</translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC đã đưc cài đt:</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation>Trò chơi đã đưc cài đt</translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG </translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>\nPKG %1\n您想要 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>\nPKG %1 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>%2\n您想要 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation></translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>\n游戏已安装%1\n您想要 </translation>
<source>PKG and Game versions match: </source>
<translation>PKG : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>%2</translation>
<source>Would you like to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>\n%1\n您想要覆盖吗</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG %1 : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC </translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation> DLC: %1 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC :</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -566,44 +566,54 @@
<translation>PKG </translation>
</message>
<message>
<location filename="../main_window.cpp" line="633"/>
<source>Patch detected!\nPKG and Game versions match!: %1\nWould you like </source>
<translation>\nPKG : %1\n您是否希望 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="634"/>
<source>to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="639"/>
<source>Patch detected!\nPKG Version %1 is older </source>
<translation>\nPKG %1 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="640"/>
<source>than installed version!: %2\nWould you like </source>
<translation>: %2\n您是否希望 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="641"/>
<source>to overwrite?</source>
<translation></translation>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="646"/>
<source>Patch detected!\nGame is installed: %1\nWould you like </source>
<translation>\n遊戲已安裝: %1\n您是否希望 </translation>
<source>PKG and Game versions match: </source>
<translation>PKG : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="647"/>
<source>to install Patch: %2?</source>
<translation>: %2</translation>
<source>Would you like to overwrite?</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="659"/>
<source>Game already installed\n%1\nWould you like to overwrite?</source>
<translation>\n%1\n您是否希望覆蓋</translation>
<location filename="../main_window.cpp" line="639"/>
<source>PKG Version %1 is older than installed version: </source>
<translation>PKG %1 : </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Game is installed: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="660"/>
<source>Would you like to install Patch: </source>
<translation>: </translation>
</message>
<message>
<location filename="../main_window.cpp" line="673"/>
<source>DLC Installation</source>
<translation>DLC </translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>
<source>Would you like to install DLC: %1?</source>
<translation> DLC: %1 </translation>
</message>
<message>
<location filename="../main_window.cpp" line="688"/>
<source>DLC already installed:</source>
<translation>DLC :</translation>
</message>
<message>
<location filename="../main_window.cpp" line="701"/>
<source>Game already installed</source>
<translation></translation>
</message>
<message>
<location filename="../main_window.cpp" line="674"/>

View File

@ -99,7 +99,7 @@ Id TypeId(const EmitContext& ctx, IR::Type type) {
}
}
void Traverse(EmitContext& ctx, IR::Program& program) {
void Traverse(EmitContext& ctx, const IR::Program& program) {
IR::Block* current_block{};
for (const IR::AbstractSyntaxNode& node : program.syntax_list) {
switch (node.type) {
@ -162,7 +162,7 @@ void Traverse(EmitContext& ctx, IR::Program& program) {
}
}
Id DefineMain(EmitContext& ctx, IR::Program& program) {
Id DefineMain(EmitContext& ctx, const IR::Program& program) {
const Id void_function{ctx.TypeFunction(ctx.void_id)};
const Id main{ctx.OpFunction(ctx.void_id, spv::FunctionControlMask::MaskNone, void_function)};
for (IR::Block* const block : program.blocks) {
@ -185,8 +185,28 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
ctx.AddCapability(spv::Capability::Int16);
}
ctx.AddCapability(spv::Capability::Int64);
if (info.has_storage_images) {
if (info.has_storage_images || info.has_image_buffers) {
ctx.AddCapability(spv::Capability::StorageImageExtendedFormats);
ctx.AddCapability(spv::Capability::StorageImageReadWithoutFormat);
ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat);
}
if (info.has_texel_buffers) {
ctx.AddCapability(spv::Capability::SampledBuffer);
}
if (info.has_image_buffers) {
ctx.AddCapability(spv::Capability::ImageBuffer);
}
if (info.has_image_gather) {
ctx.AddCapability(spv::Capability::ImageGatherExtended);
}
if (info.has_image_query) {
ctx.AddCapability(spv::Capability::ImageQuery);
}
if (info.uses_lane_id) {
ctx.AddCapability(spv::Capability::GroupNonUniform);
}
if (info.uses_group_quad) {
ctx.AddCapability(spv::Capability::GroupNonUniformQuad);
}
switch (program.info.stage) {
case Stage::Compute: {
@ -206,19 +226,9 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
} else {
ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
}
ctx.AddCapability(spv::Capability::GroupNonUniform);
if (info.uses_group_quad) {
ctx.AddCapability(spv::Capability::GroupNonUniformQuad);
}
if (info.has_discard) {
ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
}
if (info.has_image_gather) {
ctx.AddCapability(spv::Capability::ImageGatherExtended);
}
if (info.has_image_query) {
ctx.AddCapability(spv::Capability::ImageQuery);
}
if (info.stores.Get(IR::Attribute::Depth)) {
ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing);
}
@ -229,7 +239,7 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
ctx.AddEntryPoint(execution_model, main, "main", interfaces);
}
void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
void PatchPhiNodes(const IR::Program& program, EmitContext& ctx) {
auto inst{program.blocks.front()->begin()};
size_t block_index{0};
ctx.PatchDeferredPhi([&](size_t phi_arg) {
@ -248,8 +258,8 @@ void PatchPhiNodes(IR::Program& program, EmitContext& ctx) {
}
} // Anonymous namespace
std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) {
EmitContext ctx{profile, program, binding};
std::vector<u32> EmitSPIRV(const Profile& profile, const IR::Program& program, u32& binding) {
EmitContext ctx{profile, program.info, binding};
const Id main{DefineMain(ctx, program)};
DefineEntryPoint(program, ctx, main);
if (program.info.stage == Stage::Vertex) {

View File

@ -9,7 +9,7 @@
namespace Shader::Backend::SPIRV {
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, IR::Program& program,
[[nodiscard]] std::vector<u32> EmitSPIRV(const Profile& profile, const IR::Program& program,
u32& binding);
} // namespace Shader::Backend::SPIRV

View File

@ -262,171 +262,16 @@ Id EmitLoadBufferF32x4(EmitContext& ctx, IR::Inst*, u32 handle, Id address) {
return EmitLoadBufferF32xN<4>(ctx, handle, address);
}
static bool IsSignedInteger(AmdGpu::NumberFormat format) {
switch (format) {
case AmdGpu::NumberFormat::Unorm:
case AmdGpu::NumberFormat::Uscaled:
case AmdGpu::NumberFormat::Uint:
return false;
case AmdGpu::NumberFormat::Snorm:
case AmdGpu::NumberFormat::Sscaled:
case AmdGpu::NumberFormat::Sint:
case AmdGpu::NumberFormat::SnormNz:
return true;
case AmdGpu::NumberFormat::Float:
default:
UNREACHABLE();
}
}
static u32 UXBitsMax(u32 bit_width) {
return (1u << bit_width) - 1u;
}
static u32 SXBitsMax(u32 bit_width) {
return (1u << (bit_width - 1u)) - 1u;
}
static Id ConvertValue(EmitContext& ctx, Id value, AmdGpu::NumberFormat format, u32 bit_width) {
switch (format) {
case AmdGpu::NumberFormat::Unorm:
return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(UXBitsMax(bit_width))));
case AmdGpu::NumberFormat::Snorm:
return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(SXBitsMax(bit_width))));
case AmdGpu::NumberFormat::SnormNz:
// (x * 2 + 1) / (Format::SMAX * 2)
value = ctx.OpFMul(ctx.F32[1], value, ctx.ConstF32(2.f));
value = ctx.OpFAdd(ctx.F32[1], value, ctx.ConstF32(1.f));
return ctx.OpFDiv(ctx.F32[1], value, ctx.ConstF32(float(SXBitsMax(bit_width) * 2)));
case AmdGpu::NumberFormat::Uscaled:
case AmdGpu::NumberFormat::Sscaled:
case AmdGpu::NumberFormat::Uint:
case AmdGpu::NumberFormat::Sint:
case AmdGpu::NumberFormat::Float:
return value;
default:
UNREACHABLE_MSG("Unsupported number format for conversion: {}",
magic_enum::enum_name(format));
}
}
static Id ComponentOffset(EmitContext& ctx, Id address, u32 stride, u32 bit_offset) {
Id comp_offset = ctx.ConstU32(bit_offset);
if (stride < 4) {
// comp_offset += (address % 4) * 8;
const Id byte_offset = ctx.OpUMod(ctx.U32[1], address, ctx.ConstU32(4u));
const Id bit_offset = ctx.OpShiftLeftLogical(ctx.U32[1], byte_offset, ctx.ConstU32(3u));
comp_offset = ctx.OpIAdd(ctx.U32[1], comp_offset, bit_offset);
}
return comp_offset;
}
static Id GetBufferFormatValue(EmitContext& ctx, u32 handle, Id address, u32 comp) {
auto& buffer = ctx.buffers[handle];
const auto format = buffer.dfmt;
switch (format) {
case AmdGpu::DataFormat::FormatInvalid:
return ctx.f32_zero_value;
case AmdGpu::DataFormat::Format8:
case AmdGpu::DataFormat::Format16:
case AmdGpu::DataFormat::Format32:
case AmdGpu::DataFormat::Format8_8:
case AmdGpu::DataFormat::Format16_16:
case AmdGpu::DataFormat::Format10_11_11:
case AmdGpu::DataFormat::Format11_11_10:
case AmdGpu::DataFormat::Format10_10_10_2:
case AmdGpu::DataFormat::Format2_10_10_10:
case AmdGpu::DataFormat::Format8_8_8_8:
case AmdGpu::DataFormat::Format32_32:
case AmdGpu::DataFormat::Format16_16_16_16:
case AmdGpu::DataFormat::Format32_32_32:
case AmdGpu::DataFormat::Format32_32_32_32: {
const u32 num_components = AmdGpu::NumComponents(format);
if (comp >= num_components) {
return ctx.f32_zero_value;
}
// uint index = address / 4;
Id index = ctx.OpShiftRightLogical(ctx.U32[1], address, ctx.ConstU32(2u));
const u32 stride = buffer.stride;
if (stride > 4) {
const u32 index_offset = u32(AmdGpu::ComponentOffset(format, comp) / 32);
if (index_offset > 0) {
// index += index_offset;
index = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(index_offset));
}
}
const Id ptr = ctx.OpAccessChain(buffer.pointer_type, buffer.id, ctx.u32_zero_value, index);
const u32 bit_offset = AmdGpu::ComponentOffset(format, comp) % 32;
const u32 bit_width = AmdGpu::ComponentBits(format, comp);
const auto num_format = buffer.nfmt;
if (num_format == AmdGpu::NumberFormat::Float) {
if (bit_width == 32) {
return ctx.OpLoad(ctx.F32[1], ptr);
} else if (bit_width == 16) {
const Id comp_offset = ComponentOffset(ctx, address, stride, bit_offset);
Id value = ctx.OpLoad(ctx.U32[1], ptr);
value =
ctx.OpBitFieldSExtract(ctx.S32[1], value, comp_offset, ctx.ConstU32(bit_width));
value = ctx.OpSConvert(ctx.U16, value);
value = ctx.OpBitcast(ctx.F16[1], value);
return ctx.OpFConvert(ctx.F32[1], value);
} else {
UNREACHABLE_MSG("Invalid float bit width {}", bit_width);
}
} else {
Id value = ctx.OpLoad(ctx.U32[1], ptr);
const bool is_signed = IsSignedInteger(num_format);
if (bit_width < 32) {
const Id comp_offset = ComponentOffset(ctx, address, stride, bit_offset);
if (is_signed) {
value = ctx.OpBitFieldSExtract(ctx.S32[1], value, comp_offset,
ctx.ConstU32(bit_width));
} else {
value = ctx.OpBitFieldUExtract(ctx.U32[1], value, comp_offset,
ctx.ConstU32(bit_width));
}
}
value = ctx.OpBitcast(ctx.F32[1], value);
return ConvertValue(ctx, value, num_format, bit_width);
}
break;
}
default:
UNREACHABLE_MSG("Invalid format for conversion: {}", magic_enum::enum_name(format));
}
}
template <u32 N>
static Id EmitLoadBufferFormatF32xN(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
auto& buffer = ctx.buffers[handle];
address = ctx.OpIAdd(ctx.U32[1], address, buffer.offset);
if constexpr (N == 1) {
return GetBufferFormatValue(ctx, handle, address, 0);
} else {
boost::container::static_vector<Id, N> ids;
for (u32 i = 0; i < N; i++) {
ids.push_back(GetBufferFormatValue(ctx, handle, address, i));
}
return ctx.OpCompositeConstruct(ctx.F32[N], ids);
}
}
Id EmitLoadBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
return EmitLoadBufferFormatF32xN<1>(ctx, inst, handle, address);
}
Id EmitLoadBufferFormatF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
return EmitLoadBufferFormatF32xN<2>(ctx, inst, handle, address);
}
Id EmitLoadBufferFormatF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
return EmitLoadBufferFormatF32xN<3>(ctx, inst, handle, address);
}
Id EmitLoadBufferFormatF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address) {
return EmitLoadBufferFormatF32xN<4>(ctx, inst, handle, address);
const auto& buffer = ctx.texture_buffers[handle];
const Id tex_buffer = ctx.OpLoad(buffer.image_type, buffer.id);
const Id coord = ctx.OpIAdd(ctx.U32[1], address, buffer.coord_offset);
Id texel = buffer.is_storage ? ctx.OpImageRead(buffer.result_type, tex_buffer, coord)
: ctx.OpImageFetch(buffer.result_type, tex_buffer, coord);
if (buffer.is_integer) {
texel = ctx.OpBitcast(ctx.F32[4], texel);
}
return texel;
}
template <u32 N>
@ -467,97 +312,14 @@ void EmitStoreBufferU32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address
EmitStoreBufferF32xN<1>(ctx, handle, address, value);
}
static Id ConvertF32ToFormat(EmitContext& ctx, Id value, AmdGpu::NumberFormat format,
u32 bit_width) {
switch (format) {
case AmdGpu::NumberFormat::Unorm:
return ctx.OpConvertFToU(
ctx.U32[1], ctx.OpFMul(ctx.F32[1], value, ctx.ConstF32(float(UXBitsMax(bit_width)))));
case AmdGpu::NumberFormat::Uint:
return ctx.OpBitcast(ctx.U32[1], value);
case AmdGpu::NumberFormat::Float:
return value;
default:
UNREACHABLE_MSG("Unsupported number format for conversion: {}",
magic_enum::enum_name(format));
}
}
template <u32 N>
static void EmitStoreBufferFormatF32xN(EmitContext& ctx, u32 handle, Id address, Id value) {
auto& buffer = ctx.buffers[handle];
const auto format = buffer.dfmt;
const auto num_format = buffer.nfmt;
switch (format) {
case AmdGpu::DataFormat::FormatInvalid:
return;
case AmdGpu::DataFormat::Format8_8_8_8:
case AmdGpu::DataFormat::Format16:
case AmdGpu::DataFormat::Format32:
case AmdGpu::DataFormat::Format32_32:
case AmdGpu::DataFormat::Format32_32_32_32: {
ASSERT(N == AmdGpu::NumComponents(format));
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);
Id packed_value{};
for (u32 i = 0; i < N; i++) {
const u32 bit_width = AmdGpu::ComponentBits(format, i);
const u32 bit_offset = AmdGpu::ComponentOffset(format, i) % 32;
const Id comp{ConvertF32ToFormat(
ctx, N == 1 ? value : ctx.OpCompositeExtract(ctx.F32[1], value, i), num_format,
bit_width)};
if (bit_width == 32) {
if constexpr (N == 1) {
ctx.OpStore(ptr, comp);
} else {
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, comp);
}
} else {
if (i == 0) {
packed_value = comp;
} else {
packed_value =
ctx.OpBitFieldInsert(ctx.U32[1], packed_value, comp,
ctx.ConstU32(bit_offset), ctx.ConstU32(bit_width));
}
if (i == N - 1) {
ctx.OpStore(ptr, packed_value);
}
}
}
} break;
default:
UNREACHABLE_MSG("Invalid format for conversion: {}", magic_enum::enum_name(format));
}
}
void EmitStoreBufferFormatF32(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address, Id value) {
EmitStoreBufferFormatF32xN<1>(ctx, handle, address, value);
}
void EmitStoreBufferFormatF32x2(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
Id value) {
EmitStoreBufferFormatF32xN<2>(ctx, handle, address, value);
}
void EmitStoreBufferFormatF32x3(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
Id value) {
EmitStoreBufferFormatF32xN<3>(ctx, handle, address, value);
}
void EmitStoreBufferFormatF32x4(EmitContext& ctx, IR::Inst* inst, u32 handle, Id address,
Id value) {
EmitStoreBufferFormatF32xN<4>(ctx, handle, address, value);
const auto& buffer = ctx.texture_buffers[handle];
const Id tex_buffer = ctx.OpLoad(buffer.image_type, buffer.id);
const Id coord = ctx.OpIAdd(ctx.U32[1], address, buffer.coord_offset);
if (buffer.is_integer) {
value = ctx.OpBitcast(ctx.U32[4], value);
}
ctx.OpImageWrite(tex_buffer, coord, value);
}
} // namespace Shader::Backend::SPIRV

View File

@ -41,13 +41,14 @@ void Name(EmitContext& ctx, Id object, std::string_view format_str, Args&&... ar
} // Anonymous namespace
EmitContext::EmitContext(const Profile& profile_, IR::Program& program, u32& binding_)
: Sirit::Module(profile_.supported_spirv), info{program.info}, profile{profile_},
stage{program.info.stage}, binding{binding_} {
EmitContext::EmitContext(const Profile& profile_, const Shader::Info& info_, u32& binding_)
: Sirit::Module(profile_.supported_spirv), info{info_}, profile{profile_}, stage{info.stage},
binding{binding_} {
AddCapability(spv::Capability::Shader);
DefineArithmeticTypes();
DefineInterfaces();
DefineBuffers();
DefineTextureBuffers();
DefineImagesAndSamplers();
DefineSharedMemory();
}
@ -123,25 +124,24 @@ void EmitContext::DefineInterfaces() {
DefineOutputs();
}
Id GetAttributeType(EmitContext& ctx, AmdGpu::NumberFormat fmt) {
const VectorIds& GetAttributeType(EmitContext& ctx, AmdGpu::NumberFormat fmt) {
switch (fmt) {
case AmdGpu::NumberFormat::Float:
case AmdGpu::NumberFormat::Unorm:
case AmdGpu::NumberFormat::Snorm:
case AmdGpu::NumberFormat::SnormNz:
return ctx.F32[4];
case AmdGpu::NumberFormat::Sint:
return ctx.S32[4];
case AmdGpu::NumberFormat::Uint:
return ctx.U32[4];
case AmdGpu::NumberFormat::Sscaled:
return ctx.F32[4];
case AmdGpu::NumberFormat::Uscaled:
return ctx.F32[4];
case AmdGpu::NumberFormat::Srgb:
return ctx.F32;
case AmdGpu::NumberFormat::Sint:
return ctx.S32;
case AmdGpu::NumberFormat::Uint:
return ctx.U32;
default:
break;
}
throw InvalidArgument("Invalid attribute type {}", fmt);
UNREACHABLE_MSG("Invalid attribute type {}", fmt);
}
EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat fmt, Id id) {
@ -162,7 +162,7 @@ EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat f
default:
break;
}
throw InvalidArgument("Invalid attribute type {}", fmt);
UNREACHABLE_MSG("Invalid attribute type {}", fmt);
}
void EmitContext::DefineBufferOffsets() {
@ -177,6 +177,16 @@ void EmitContext::DefineBufferOffsets() {
buffer.offset = OpBitFieldUExtract(U32[1], value, ConstU32(offset), ConstU32(8U));
buffer.offset_dwords = OpShiftRightLogical(U32[1], buffer.offset, ConstU32(2U));
}
for (auto& tex_buffer : texture_buffers) {
const u32 binding = tex_buffer.binding;
const u32 half = Shader::PushData::BufOffsetIndex + (binding >> 4);
const u32 comp = (binding & 0xf) >> 2;
const u32 offset = (binding & 0x3) << 3;
const Id ptr{OpAccessChain(TypePointer(spv::StorageClass::PushConstant, U32[1]),
push_data_block, ConstU32(half), ConstU32(comp))};
const Id value{OpLoad(U32[1], ptr)};
tex_buffer.coord_offset = OpBitFieldUExtract(U32[1], value, ConstU32(offset), ConstU32(8U));
}
}
Id MakeDefaultValue(EmitContext& ctx, u32 default_value) {
@ -195,6 +205,11 @@ Id MakeDefaultValue(EmitContext& ctx, u32 default_value) {
}
void EmitContext::DefineInputs() {
if (info.uses_lane_id) {
subgroup_local_invocation_id = DefineVariable(
U32[1], spv::BuiltIn::SubgroupLocalInvocationId, spv::StorageClass::Input);
Decorate(subgroup_local_invocation_id, spv::Decoration::Flat);
}
switch (stage) {
case Stage::Vertex: {
vertex_index = DefineVariable(U32[1], spv::BuiltIn::VertexIndex, spv::StorageClass::Input);
@ -202,7 +217,7 @@ void EmitContext::DefineInputs() {
instance_id = DefineVariable(U32[1], spv::BuiltIn::InstanceIndex, spv::StorageClass::Input);
for (const auto& input : info.vs_inputs) {
const Id type{GetAttributeType(*this, input.fmt)};
const Id type{GetAttributeType(*this, input.fmt)[4]};
if (input.instance_step_rate == Info::VsInput::InstanceIdType::OverStepRate0 ||
input.instance_step_rate == Info::VsInput::InstanceIdType::OverStepRate1) {
@ -229,15 +244,12 @@ void EmitContext::DefineInputs() {
break;
}
case Stage::Fragment:
subgroup_local_invocation_id = DefineVariable(
U32[1], spv::BuiltIn::SubgroupLocalInvocationId, spv::StorageClass::Input);
Decorate(subgroup_local_invocation_id, spv::Decoration::Flat);
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);
for (const auto& input : info.ps_inputs) {
const u32 semantic = input.param_index;
if (input.is_default) {
if (input.is_default && !input.is_flat) {
input_params[semantic] = {MakeDefaultValue(*this, input.default_value), F32[1],
F32[1], 4, true};
continue;
@ -328,47 +340,75 @@ void EmitContext::DefinePushDataBlock() {
void EmitContext::DefineBuffers() {
boost::container::small_vector<Id, 8> type_ids;
for (u32 i = 0; const auto& buffer : info.buffers) {
const auto* data_types = True(buffer.used_types & IR::Type::F32) ? &F32 : &U32;
const Id data_type = (*data_types)[1];
const Id record_array_type{buffer.is_storage
? TypeRuntimeArray(data_type)
: TypeArray(data_type, ConstU32(buffer.length))};
const auto define_struct = [&](Id record_array_type, bool is_instance_data) {
const Id struct_type{TypeStruct(record_array_type)};
if (std::ranges::find(type_ids, record_array_type.value, &Id::value) == type_ids.end()) {
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);
const auto name =
buffer.is_instance_data
? fmt::format("{}_instance_data{}_{}{}", stage, i, 'f',
sizeof(float) * CHAR_BIT)
: fmt::format("{}_cbuf_block_{}{}", stage, 'f', sizeof(float) * CHAR_BIT);
const auto name = is_instance_data ? fmt::format("{}_instance_data_f32", stage)
: fmt::format("{}_cbuf_block_f32", stage);
Name(struct_type, name);
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;
};
for (const auto& desc : info.buffers) {
const auto sharp = desc.GetSharp(info);
const bool is_storage = desc.IsStorage(sharp);
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(sharp.NumDwords()))};
const Id struct_type{define_struct(record_array_type, desc.is_instance_data)};
const auto storage_class =
buffer.is_storage ? spv::StorageClass::StorageBuffer : spv::StorageClass::Uniform;
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);
Decorate(id, spv::Decoration::DescriptorSet, 0U);
Name(id, fmt::format("{}_{}", buffer.is_storage ? "ssbo" : "cbuf", buffer.sgpr_base));
if (is_storage && !desc.is_written) {
Decorate(id, spv::Decoration::NonWritable);
}
Name(id, fmt::format("{}_{}", is_storage ? "ssbo" : "cbuf", desc.sgpr_base));
buffers.push_back({
.id = id,
.binding = binding++,
.data_types = data_types,
.pointer_type = pointer_type,
.dfmt = buffer.dfmt,
.nfmt = buffer.nfmt,
.stride = buffer.GetVsharp(info).GetStride(),
});
interfaces.push_back(id);
i++;
}
}
void EmitContext::DefineTextureBuffers() {
for (const auto& desc : info.texture_buffers) {
const bool is_integer =
desc.nfmt == AmdGpu::NumberFormat::Uint || desc.nfmt == AmdGpu::NumberFormat::Sint;
const VectorIds& sampled_type{GetAttributeType(*this, desc.nfmt)};
const u32 sampled = desc.is_written ? 2 : 1;
const Id image_type{TypeImage(sampled_type[1], spv::Dim::Buffer, false, false, false,
sampled, spv::ImageFormat::Unknown)};
const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)};
const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)};
Decorate(id, spv::Decoration::Binding, binding);
Decorate(id, spv::Decoration::DescriptorSet, 0U);
Name(id, fmt::format("{}_{}", desc.is_written ? "imgbuf" : "texbuf", desc.sgpr_base));
texture_buffers.push_back({
.id = id,
.binding = binding++,
.image_type = image_type,
.result_type = sampled_type[4],
.is_integer = is_integer,
.is_storage = desc.is_written,
});
interfaces.push_back(id);
}
}
@ -447,7 +487,7 @@ spv::ImageFormat GetFormat(const AmdGpu::Image& image) {
Id ImageType(EmitContext& ctx, const ImageResource& desc, Id sampled_type) {
const auto image = ctx.info.ReadUd<AmdGpu::Image>(desc.sgpr_base, desc.dword_offset);
const auto format = desc.is_storage ? GetFormat(image) : spv::ImageFormat::Unknown;
const auto format = desc.is_atomic ? GetFormat(image) : spv::ImageFormat::Unknown;
const u32 sampled = desc.is_storage ? 2 : 1;
switch (desc.type) {
case AmdGpu::ImageType::Color1D:
@ -470,17 +510,8 @@ Id ImageType(EmitContext& ctx, const ImageResource& desc, Id sampled_type) {
void EmitContext::DefineImagesAndSamplers() {
for (const auto& image_desc : info.images) {
const VectorIds* data_types = [&] {
switch (image_desc.nfmt) {
case AmdGpu::NumberFormat::Uint:
return &U32;
case AmdGpu::NumberFormat::Sint:
return &S32;
default:
return &F32;
}
}();
const Id sampled_type = data_types->Get(1);
const VectorIds& data_types = GetAttributeType(*this, image_desc.nfmt);
const Id sampled_type = data_types[1];
const Id image_type{ImageType(*this, image_desc, sampled_type)};
const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)};
const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)};
@ -489,7 +520,7 @@ void EmitContext::DefineImagesAndSamplers() {
Name(id, fmt::format("{}_{}{}_{:02x}", stage, "img", image_desc.sgpr_base,
image_desc.dword_offset));
images.push_back({
.data_types = data_types,
.data_types = &data_types,
.id = id,
.sampled_type = image_desc.is_storage ? sampled_type : TypeSampledImage(image_type),
.pointer_type = pointer_type,
@ -498,13 +529,12 @@ void EmitContext::DefineImagesAndSamplers() {
interfaces.push_back(id);
++binding;
}
if (std::ranges::any_of(info.images, &ImageResource::is_atomic)) {
image_u32 = TypePointer(spv::StorageClass::Image, U32[1]);
}
if (info.samplers.empty()) {
return;
}
sampler_type = TypeSampler();
sampler_pointer_type = TypePointer(spv::StorageClass::UniformConstant, sampler_type);
for (const auto& samp_desc : info.samplers) {
@ -520,14 +550,15 @@ void EmitContext::DefineImagesAndSamplers() {
}
void EmitContext::DefineSharedMemory() {
static constexpr size_t DefaultSharedMemSize = 16_KB;
static constexpr size_t DefaultSharedMemSize = 2_KB;
if (!info.uses_shared) {
return;
}
if (info.shared_memory_size == 0) {
info.shared_memory_size = DefaultSharedMemSize;
u32 shared_memory_size = info.shared_memory_size;
if (shared_memory_size == 0) {
shared_memory_size = DefaultSharedMemSize;
}
const u32 num_elements{Common::DivCeil(info.shared_memory_size, 4U)};
const u32 num_elements{Common::DivCeil(shared_memory_size, 4U)};
const Id type{TypeArray(U32[1], ConstU32(num_elements))};
shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type);
shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]);

View File

@ -36,7 +36,7 @@ struct VectorIds {
class EmitContext final : public Sirit::Module {
public:
explicit EmitContext(const Profile& profile, IR::Program& program, u32& binding);
explicit EmitContext(const Profile& profile, const Shader::Info& info, u32& binding);
~EmitContext();
Id Def(const IR::Value& value);
@ -124,7 +124,7 @@ public:
return ConstantComposite(type, constituents);
}
Info& info;
const Info& info;
const Profile& profile;
Stage stage{};
@ -207,13 +207,20 @@ public:
u32 binding;
const VectorIds* data_types;
Id pointer_type;
AmdGpu::DataFormat dfmt;
AmdGpu::NumberFormat nfmt;
u32 stride;
};
struct TextureBufferDefinition {
Id id;
Id coord_offset;
u32 binding;
Id image_type;
Id result_type;
bool is_integer;
bool is_storage;
};
u32& binding;
boost::container::small_vector<BufferDefinition, 16> buffers;
boost::container::small_vector<TextureBufferDefinition, 8> texture_buffers;
boost::container::small_vector<TextureDefinition, 8> images;
boost::container::small_vector<Id, 4> samplers;
@ -238,6 +245,7 @@ private:
void DefineOutputs();
void DefinePushDataBlock();
void DefineBuffers();
void DefineTextureBuffers();
void DefineImagesAndSamplers();
void DefineSharedMemory();

View File

@ -18,25 +18,31 @@ void Translator::EmitDataShare(const GcnInst& inst) {
case Opcode::DS_READ2_B64:
return DS_READ(64, false, true, inst);
case Opcode::DS_WRITE_B32:
return DS_WRITE(32, false, false, inst);
return DS_WRITE(32, false, false, false, inst);
case Opcode::DS_WRITE2ST64_B32:
return DS_WRITE(32, false, true, true, inst);
case Opcode::DS_WRITE_B64:
return DS_WRITE(64, false, false, inst);
return DS_WRITE(64, false, false, false, inst);
case Opcode::DS_WRITE2_B32:
return DS_WRITE(32, false, true, inst);
return DS_WRITE(32, false, true, false, inst);
case Opcode::DS_WRITE2_B64:
return DS_WRITE(64, false, true, inst);
return DS_WRITE(64, false, true, false, inst);
case Opcode::DS_ADD_U32:
return DS_ADD_U32(inst, false);
case Opcode::DS_MIN_U32:
return DS_MIN_U32(inst, false);
return DS_MIN_U32(inst, false, false);
case Opcode::DS_MIN_I32:
return DS_MIN_U32(inst, true, false);
case Opcode::DS_MAX_U32:
return DS_MAX_U32(inst, false);
return DS_MAX_U32(inst, false, false);
case Opcode::DS_MAX_I32:
return DS_MAX_U32(inst, true, false);
case Opcode::DS_ADD_RTN_U32:
return DS_ADD_U32(inst, true);
case Opcode::DS_MIN_RTN_U32:
return DS_MIN_U32(inst, true);
return DS_MIN_U32(inst, false, true);
case Opcode::DS_MAX_RTN_U32:
return DS_MAX_U32(inst, true);
return DS_MAX_U32(inst, false, true);
default:
LogMissingOpcode(inst);
}
@ -89,12 +95,13 @@ void Translator::DS_READ(int bit_size, bool is_signed, bool is_pair, const GcnIn
}
}
void Translator::DS_WRITE(int bit_size, bool is_signed, bool is_pair, const GcnInst& inst) {
void Translator::DS_WRITE(int bit_size, bool is_signed, bool is_pair, bool stride64,
const GcnInst& inst) {
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};
if (is_pair) {
const u32 adj = bit_size == 32 ? 4 : 8;
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)));
if (bit_size == 32) {
ir.WriteShared(32, ir.GetVectorReg(data0), addr0);
@ -133,23 +140,23 @@ void Translator::DS_ADD_U32(const GcnInst& inst, bool rtn) {
}
}
void Translator::DS_MIN_U32(const GcnInst& inst, bool rtn) {
void Translator::DS_MIN_U32(const GcnInst& inst, bool is_signed, bool rtn) {
const IR::U32 addr{GetSrc(inst.src[0])};
const IR::U32 data{GetSrc(inst.src[1])};
const IR::U32 offset = ir.Imm32(u32(inst.control.ds.offset0));
const IR::U32 addr_offset = ir.IAdd(addr, offset);
const IR::Value original_val = ir.SharedAtomicIMin(addr_offset, data, false);
const IR::Value original_val = ir.SharedAtomicIMin(addr_offset, data, is_signed);
if (rtn) {
SetDst(inst.dst[0], IR::U32{original_val});
}
}
void Translator::DS_MAX_U32(const GcnInst& inst, bool rtn) {
void Translator::DS_MAX_U32(const GcnInst& inst, bool is_signed, bool rtn) {
const IR::U32 addr{GetSrc(inst.src[0])};
const IR::U32 data{GetSrc(inst.src[1])};
const IR::U32 offset = ir.Imm32(u32(inst.control.ds.offset0));
const IR::U32 addr_offset = ir.IAdd(addr, offset);
const IR::Value original_val = ir.SharedAtomicIMax(addr_offset, data, false);
const IR::Value original_val = ir.SharedAtomicIMax(addr_offset, data, is_signed);
if (rtn) {
SetDst(inst.dst[0], IR::U32{original_val});
}

View File

@ -1,14 +1,12 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "common/logging/log.h"
#include "shader_recompiler/frontend/translate/translate.h"
namespace Shader::Gcn {
void Translator::EmitExport(const GcnInst& inst) {
if (ir.block->has_multiple_predecessors && info.stage == Stage::Fragment) {
LOG_WARNING(Render_Recompiler, "An ambiguous export appeared in translation");
ir.Discard(ir.LogicalNot(ir.GetExec()));
}

View File

@ -31,6 +31,8 @@ void Translator::EmitScalarAlu(const GcnInst& inst) {
return S_OR_B64(NegateMode::Result, false, inst);
case Opcode::S_XOR_B64:
return S_OR_B64(NegateMode::None, true, inst);
case Opcode::S_XNOR_B64:
return S_OR_B64(NegateMode::Result, true, inst);
case Opcode::S_ORN2_B64:
return S_OR_B64(NegateMode::Src1, false, inst);
case Opcode::S_AND_B64:

View File

@ -354,7 +354,7 @@ void Translator::EmitFetch(const GcnInst& inst) {
if (!std::filesystem::exists(dump_dir)) {
std::filesystem::create_directories(dump_dir);
}
const auto filename = fmt::format("vs_fetch_{:#018x}.bin", info.pgm_hash);
const auto filename = fmt::format("vs_{:#018x}_fetch.bin", info.pgm_hash);
const auto file = IOFile{dump_dir / filename, FileAccessMode::Write};
file.WriteRaw<u8>(code, fetch_size);
}
@ -399,9 +399,7 @@ void Translator::EmitFetch(const GcnInst& inst) {
info.buffers.push_back({
.sgpr_base = attrib.sgpr_base,
.dword_offset = attrib.dword_offset,
.length = buffer.num_records,
.used_types = IR::Type::F32,
.is_storage = true, // we may not fit into UBO with large meshes
.is_instance_data = true,
});
instance_buf_handle = s32(info.buffers.size() - 1);
@ -438,6 +436,7 @@ void Translator::EmitFlowControl(u32 pc, const GcnInst& inst) {
case Opcode::S_CBRANCH_SCC1:
case Opcode::S_CBRANCH_VCCNZ:
case Opcode::S_CBRANCH_VCCZ:
case Opcode::S_CBRANCH_EXECNZ:
case Opcode::S_BRANCH:
return;
default:

View File

@ -191,8 +191,10 @@ public:
void V_MBCNT_U32_B32(bool is_low, const GcnInst& inst);
// Vector Memory
void BUFFER_LOAD_FORMAT(u32 num_dwords, bool is_typed, bool is_format, const GcnInst& inst);
void BUFFER_STORE_FORMAT(u32 num_dwords, bool is_typed, bool is_format, const GcnInst& inst);
void BUFFER_LOAD(u32 num_dwords, bool is_typed, const GcnInst& inst);
void BUFFER_LOAD_FORMAT(u32 num_dwords, const GcnInst& inst);
void BUFFER_STORE(u32 num_dwords, bool is_typed, const GcnInst& inst);
void BUFFER_STORE_FORMAT(u32 num_dwords, const GcnInst& inst);
void BUFFER_ATOMIC(AtomicOp op, const GcnInst& inst);
// Vector interpolation
@ -202,10 +204,10 @@ public:
// Data share
void DS_SWIZZLE_B32(const GcnInst& inst);
void DS_READ(int bit_size, bool is_signed, bool is_pair, const GcnInst& inst);
void DS_WRITE(int bit_size, bool is_signed, bool is_pair, const GcnInst& inst);
void DS_WRITE(int bit_size, bool is_signed, bool is_pair, bool stride64, const GcnInst& inst);
void DS_ADD_U32(const GcnInst& inst, bool rtn);
void DS_MIN_U32(const GcnInst& inst, bool rtn);
void DS_MAX_U32(const GcnInst& inst, bool rtn);
void DS_MIN_U32(const GcnInst& inst, bool is_signed, bool rtn);
void DS_MAX_U32(const GcnInst& inst, bool is_signed, bool rtn);
void V_READFIRSTLANE_B32(const GcnInst& inst);
void V_READLANE_B32(const GcnInst& inst);
void V_WRITELANE_B32(const GcnInst& inst);

View File

@ -415,14 +415,20 @@ void Translator::V_ADDC_U32(const GcnInst& inst) {
const auto src0 = GetSrc<IR::U32>(inst.src[0]);
const auto src1 = GetSrc<IR::U32>(inst.src[1]);
IR::U32 scarry;
IR::U1 carry;
if (inst.src_count == 3) { // VOP3
IR::U1 thread_bit{ir.GetThreadBitScalarReg(IR::ScalarReg(inst.src[2].code))};
scarry = IR::U32{ir.Select(thread_bit, ir.Imm32(1), ir.Imm32(0))};
if (inst.src[2].field == OperandField::VccLo) {
carry = ir.GetVcc();
} else if (inst.src[2].field == OperandField::ScalarGPR) {
carry = ir.GetThreadBitScalarReg(IR::ScalarReg(inst.src[2].code));
} else {
UNREACHABLE();
}
} else { // VOP2
scarry = ir.GetVccLo();
carry = ir.GetVcc();
}
const IR::U32 scarry = IR::U32{ir.Select(carry, ir.Imm32(1), ir.Imm32(0))};
const IR::U32 result = ir.IAdd(ir.IAdd(src0, src1), scarry);
const IR::VectorReg dst_reg{inst.dst[0].code};

View File

@ -56,57 +56,57 @@ void Translator::EmitVectorMemory(const GcnInst& inst) {
// Buffer load operations
case Opcode::TBUFFER_LOAD_FORMAT_X:
return BUFFER_LOAD_FORMAT(1, true, true, inst);
return BUFFER_LOAD(1, true, inst);
case Opcode::TBUFFER_LOAD_FORMAT_XY:
return BUFFER_LOAD_FORMAT(2, true, true, inst);
return BUFFER_LOAD(2, true, inst);
case Opcode::TBUFFER_LOAD_FORMAT_XYZ:
return BUFFER_LOAD_FORMAT(3, true, true, inst);
return BUFFER_LOAD(3, true, inst);
case Opcode::TBUFFER_LOAD_FORMAT_XYZW:
return BUFFER_LOAD_FORMAT(4, true, true, inst);
return BUFFER_LOAD(4, true, inst);
case Opcode::BUFFER_LOAD_FORMAT_X:
return BUFFER_LOAD_FORMAT(1, false, true, inst);
return BUFFER_LOAD_FORMAT(1, inst);
case Opcode::BUFFER_LOAD_FORMAT_XY:
return BUFFER_LOAD_FORMAT(2, false, true, inst);
return BUFFER_LOAD_FORMAT(2, inst);
case Opcode::BUFFER_LOAD_FORMAT_XYZ:
return BUFFER_LOAD_FORMAT(3, false, true, inst);
return BUFFER_LOAD_FORMAT(3, inst);
case Opcode::BUFFER_LOAD_FORMAT_XYZW:
return BUFFER_LOAD_FORMAT(4, false, true, inst);
return BUFFER_LOAD_FORMAT(4, inst);
case Opcode::BUFFER_LOAD_DWORD:
return BUFFER_LOAD_FORMAT(1, false, false, inst);
return BUFFER_LOAD(1, false, inst);
case Opcode::BUFFER_LOAD_DWORDX2:
return BUFFER_LOAD_FORMAT(2, false, false, inst);
return BUFFER_LOAD(2, false, inst);
case Opcode::BUFFER_LOAD_DWORDX3:
return BUFFER_LOAD_FORMAT(3, false, false, inst);
return BUFFER_LOAD(3, false, inst);
case Opcode::BUFFER_LOAD_DWORDX4:
return BUFFER_LOAD_FORMAT(4, false, false, inst);
return BUFFER_LOAD(4, false, inst);
// Buffer store operations
case Opcode::BUFFER_STORE_FORMAT_X:
return BUFFER_STORE_FORMAT(1, false, true, inst);
return BUFFER_STORE_FORMAT(1, inst);
case Opcode::BUFFER_STORE_FORMAT_XY:
return BUFFER_STORE_FORMAT(2, false, true, inst);
return BUFFER_STORE_FORMAT(2, inst);
case Opcode::BUFFER_STORE_FORMAT_XYZ:
return BUFFER_STORE_FORMAT(3, false, true, inst);
return BUFFER_STORE_FORMAT(3, inst);
case Opcode::BUFFER_STORE_FORMAT_XYZW:
return BUFFER_STORE_FORMAT(4, false, true, inst);
return BUFFER_STORE_FORMAT(4, inst);
case Opcode::TBUFFER_STORE_FORMAT_X:
return BUFFER_STORE_FORMAT(1, true, true, inst);
return BUFFER_STORE(1, true, inst);
case Opcode::TBUFFER_STORE_FORMAT_XY:
return BUFFER_STORE_FORMAT(2, true, true, inst);
return BUFFER_STORE(2, true, inst);
case Opcode::TBUFFER_STORE_FORMAT_XYZ:
return BUFFER_STORE_FORMAT(3, true, true, inst);
return BUFFER_STORE(3, true, inst);
case Opcode::BUFFER_STORE_DWORD:
return BUFFER_STORE_FORMAT(1, false, false, inst);
return BUFFER_STORE(1, false, inst);
case Opcode::BUFFER_STORE_DWORDX2:
return BUFFER_STORE_FORMAT(2, false, false, inst);
return BUFFER_STORE(2, false, inst);
case Opcode::BUFFER_STORE_DWORDX3:
return BUFFER_STORE_FORMAT(3, false, false, inst);
return BUFFER_STORE(3, false, inst);
case Opcode::BUFFER_STORE_DWORDX4:
return BUFFER_STORE_FORMAT(4, false, false, inst);
return BUFFER_STORE(4, false, inst);
// Buffer atomic operations
case Opcode::BUFFER_ATOMIC_ADD:
@ -349,8 +349,7 @@ void Translator::IMAGE_STORE(const GcnInst& inst) {
ir.ImageWrite(handle, body, value, {});
}
void Translator::BUFFER_LOAD_FORMAT(u32 num_dwords, bool is_typed, bool is_format,
const GcnInst& inst) {
void Translator::BUFFER_LOAD(u32 num_dwords, bool is_typed, const GcnInst& inst) {
const auto& mtbuf = inst.control.mtbuf;
const IR::VectorReg vaddr{inst.src[0].code};
const IR::ScalarReg sharp{inst.src[2].code * 4};
@ -370,22 +369,19 @@ void Translator::BUFFER_LOAD_FORMAT(u32 num_dwords, bool is_typed, bool is_forma
info.index_enable.Assign(mtbuf.idxen);
info.offset_enable.Assign(mtbuf.offen);
info.inst_offset.Assign(mtbuf.offset);
info.is_typed.Assign(is_typed);
if (is_typed) {
info.dmft.Assign(static_cast<AmdGpu::DataFormat>(mtbuf.dfmt));
info.nfmt.Assign(static_cast<AmdGpu::NumberFormat>(mtbuf.nfmt));
ASSERT(info.nfmt == AmdGpu::NumberFormat::Float &&
(info.dmft == AmdGpu::DataFormat::Format32_32_32_32 ||
info.dmft == AmdGpu::DataFormat::Format32_32_32 ||
info.dmft == AmdGpu::DataFormat::Format32_32 ||
info.dmft == AmdGpu::DataFormat::Format32));
const auto dmft = static_cast<AmdGpu::DataFormat>(mtbuf.dfmt);
const auto nfmt = static_cast<AmdGpu::NumberFormat>(mtbuf.nfmt);
ASSERT(nfmt == AmdGpu::NumberFormat::Float &&
(dmft == AmdGpu::DataFormat::Format32_32_32_32 ||
dmft == AmdGpu::DataFormat::Format32_32_32 ||
dmft == AmdGpu::DataFormat::Format32_32 || dmft == AmdGpu::DataFormat::Format32));
}
const IR::Value handle =
ir.CompositeConstruct(ir.GetScalarReg(sharp), ir.GetScalarReg(sharp + 1),
ir.GetScalarReg(sharp + 2), ir.GetScalarReg(sharp + 3));
const IR::Value value = is_format ? ir.LoadBufferFormat(num_dwords, handle, address, info)
: ir.LoadBuffer(num_dwords, handle, address, info);
const IR::Value value = ir.LoadBuffer(num_dwords, handle, address, info);
const IR::VectorReg dst_reg{inst.src[1].code};
if (num_dwords == 1) {
ir.SetVectorReg(dst_reg, IR::F32{value});
@ -396,8 +392,34 @@ void Translator::BUFFER_LOAD_FORMAT(u32 num_dwords, bool is_typed, bool is_forma
}
}
void Translator::BUFFER_STORE_FORMAT(u32 num_dwords, bool is_typed, bool is_format,
const GcnInst& inst) {
void Translator::BUFFER_LOAD_FORMAT(u32 num_dwords, const GcnInst& inst) {
const auto& mubuf = inst.control.mubuf;
const IR::VectorReg vaddr{inst.src[0].code};
const IR::ScalarReg sharp{inst.src[2].code * 4};
ASSERT_MSG(!mubuf.offen && mubuf.offset == 0, "Offsets for image buffers are not supported");
const IR::Value address = [&] -> IR::Value {
if (mubuf.idxen) {
return ir.GetVectorReg(vaddr);
}
return {};
}();
const IR::Value soffset{GetSrc(inst.src[3])};
ASSERT_MSG(soffset.IsImmediate() && soffset.U32() == 0, "Non immediate offset not supported");
IR::BufferInstInfo info{};
info.index_enable.Assign(mubuf.idxen);
const IR::Value handle =
ir.CompositeConstruct(ir.GetScalarReg(sharp), ir.GetScalarReg(sharp + 1),
ir.GetScalarReg(sharp + 2), ir.GetScalarReg(sharp + 3));
const IR::Value value = ir.LoadBufferFormat(handle, address, info);
const IR::VectorReg dst_reg{inst.src[1].code};
for (u32 i = 0; i < num_dwords; i++) {
ir.SetVectorReg(dst_reg + i, IR::F32{ir.CompositeExtract(value, i)});
}
}
void Translator::BUFFER_STORE(u32 num_dwords, bool is_typed, const GcnInst& inst) {
const auto& mtbuf = inst.control.mtbuf;
const IR::VectorReg vaddr{inst.src[0].code};
const IR::ScalarReg sharp{inst.src[2].code * 4};
@ -417,45 +439,76 @@ void Translator::BUFFER_STORE_FORMAT(u32 num_dwords, bool is_typed, bool is_form
info.index_enable.Assign(mtbuf.idxen);
info.offset_enable.Assign(mtbuf.offen);
info.inst_offset.Assign(mtbuf.offset);
info.is_typed.Assign(is_typed);
if (is_typed) {
info.dmft.Assign(static_cast<AmdGpu::DataFormat>(mtbuf.dfmt));
info.nfmt.Assign(static_cast<AmdGpu::NumberFormat>(mtbuf.nfmt));
const auto dmft = static_cast<AmdGpu::DataFormat>(mtbuf.dfmt);
const auto nfmt = static_cast<AmdGpu::NumberFormat>(mtbuf.nfmt);
ASSERT(nfmt == AmdGpu::NumberFormat::Float &&
(dmft == AmdGpu::DataFormat::Format32_32_32_32 ||
dmft == AmdGpu::DataFormat::Format32_32_32 ||
dmft == AmdGpu::DataFormat::Format32_32 || dmft == AmdGpu::DataFormat::Format32));
}
IR::Value value{};
const IR::VectorReg src_reg{inst.src[1].code};
switch (num_dwords) {
case 1:
value = ir.GetVectorReg<Shader::IR::F32>(src_reg);
value = ir.GetVectorReg<IR::F32>(src_reg);
break;
case 2:
value = ir.CompositeConstruct(ir.GetVectorReg<Shader::IR::F32>(src_reg),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 1));
value = ir.CompositeConstruct(ir.GetVectorReg<IR::F32>(src_reg),
ir.GetVectorReg<IR::F32>(src_reg + 1));
break;
case 3:
value = ir.CompositeConstruct(ir.GetVectorReg<Shader::IR::F32>(src_reg),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 1),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 2));
value = ir.CompositeConstruct(ir.GetVectorReg<IR::F32>(src_reg),
ir.GetVectorReg<IR::F32>(src_reg + 1),
ir.GetVectorReg<IR::F32>(src_reg + 2));
break;
case 4:
value = ir.CompositeConstruct(ir.GetVectorReg<Shader::IR::F32>(src_reg),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 1),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 2),
ir.GetVectorReg<Shader::IR::F32>(src_reg + 3));
value = ir.CompositeConstruct(
ir.GetVectorReg<IR::F32>(src_reg), ir.GetVectorReg<IR::F32>(src_reg + 1),
ir.GetVectorReg<IR::F32>(src_reg + 2), ir.GetVectorReg<IR::F32>(src_reg + 3));
break;
}
const IR::Value handle =
ir.CompositeConstruct(ir.GetScalarReg(sharp), ir.GetScalarReg(sharp + 1),
ir.GetScalarReg(sharp + 2), ir.GetScalarReg(sharp + 3));
if (is_format) {
ir.StoreBufferFormat(num_dwords, handle, address, value, info);
} else {
ir.StoreBuffer(num_dwords, handle, address, value, info);
}
}
// TODO: U64
void Translator::BUFFER_STORE_FORMAT(u32 num_dwords, const GcnInst& inst) {
const auto& mubuf = inst.control.mubuf;
const IR::VectorReg vaddr{inst.src[0].code};
const IR::ScalarReg sharp{inst.src[2].code * 4};
ASSERT_MSG(!mubuf.offen && mubuf.offset == 0, "Offsets for image buffers are not supported");
const IR::Value address = [&] -> IR::Value {
if (mubuf.idxen) {
return ir.GetVectorReg(vaddr);
}
return {};
}();
const IR::Value soffset{GetSrc(inst.src[3])};
ASSERT_MSG(soffset.IsImmediate() && soffset.U32() == 0, "Non immediate offset not supported");
IR::BufferInstInfo info{};
info.index_enable.Assign(mubuf.idxen);
const IR::VectorReg src_reg{inst.src[1].code};
std::array<IR::Value, 4> comps{};
for (u32 i = 0; i < num_dwords; i++) {
comps[i] = ir.GetVectorReg<IR::F32>(src_reg + i);
}
for (u32 i = num_dwords; i < 4; i++) {
comps[i] = ir.Imm32(0.f);
}
const IR::Value value = ir.CompositeConstruct(comps[0], comps[1], comps[2], comps[3]);
const IR::Value handle =
ir.CompositeConstruct(ir.GetScalarReg(sharp), ir.GetScalarReg(sharp + 1),
ir.GetScalarReg(sharp + 2), ir.GetScalarReg(sharp + 3));
ir.StoreBufferFormat(handle, address, value, info);
}
void Translator::BUFFER_ATOMIC(AtomicOp op, const GcnInst& inst) {
const auto& mubuf = inst.control.mubuf;
const IR::VectorReg vaddr{inst.src[0].code};

View File

@ -325,20 +325,8 @@ Value IREmitter::LoadBuffer(int num_dwords, const Value& handle, const Value& ad
}
}
Value IREmitter::LoadBufferFormat(int num_dwords, const Value& handle, const Value& address,
BufferInstInfo info) {
switch (num_dwords) {
case 1:
Value IREmitter::LoadBufferFormat(const Value& handle, const Value& address, BufferInstInfo info) {
return Inst(Opcode::LoadBufferFormatF32, Flags{info}, handle, address);
case 2:
return Inst(Opcode::LoadBufferFormatF32x2, Flags{info}, handle, address);
case 3:
return Inst(Opcode::LoadBufferFormatF32x3, Flags{info}, handle, address);
case 4:
return Inst(Opcode::LoadBufferFormatF32x4, Flags{info}, handle, address);
default:
UNREACHABLE_MSG("Invalid number of dwords {}", num_dwords);
}
}
void IREmitter::StoreBuffer(int num_dwords, const Value& handle, const Value& address,
@ -409,24 +397,9 @@ Value IREmitter::BufferAtomicSwap(const Value& handle, const Value& address, con
return Inst(Opcode::BufferAtomicSwap32, Flags{info}, handle, address, value);
}
void IREmitter::StoreBufferFormat(int num_dwords, const Value& handle, const Value& address,
const Value& data, BufferInstInfo info) {
switch (num_dwords) {
case 1:
void IREmitter::StoreBufferFormat(const Value& handle, const Value& address, const Value& data,
BufferInstInfo info) {
Inst(Opcode::StoreBufferFormatF32, Flags{info}, handle, address, data);
break;
case 2:
Inst(Opcode::StoreBufferFormatF32x2, Flags{info}, handle, address, data);
break;
case 3:
Inst(Opcode::StoreBufferFormatF32x3, Flags{info}, handle, address, data);
break;
case 4:
Inst(Opcode::StoreBufferFormatF32x4, Flags{info}, handle, address, data);
break;
default:
UNREACHABLE_MSG("Invalid number of dwords {}", num_dwords);
}
}
U32 IREmitter::LaneId() {

View File

@ -92,12 +92,12 @@ public:
[[nodiscard]] Value LoadBuffer(int num_dwords, const Value& handle, const Value& address,
BufferInstInfo info);
[[nodiscard]] Value LoadBufferFormat(int num_dwords, const Value& handle, const Value& address,
[[nodiscard]] Value LoadBufferFormat(const Value& handle, const Value& address,
BufferInstInfo info);
void StoreBuffer(int num_dwords, const Value& handle, const Value& address, const Value& data,
BufferInstInfo info);
void StoreBufferFormat(int num_dwords, const Value& handle, const Value& address,
const Value& data, BufferInstInfo info);
void StoreBufferFormat(const Value& handle, const Value& address, const Value& data,
BufferInstInfo info);
[[nodiscard]] Value BufferAtomicIAdd(const Value& handle, const Value& address,
const Value& value, BufferInstInfo info);

View File

@ -56,9 +56,6 @@ bool Inst::MayHaveSideEffects() const noexcept {
case Opcode::StoreBufferF32x3:
case Opcode::StoreBufferF32x4:
case Opcode::StoreBufferFormatF32:
case Opcode::StoreBufferFormatF32x2:
case Opcode::StoreBufferFormatF32x3:
case Opcode::StoreBufferFormatF32x4:
case Opcode::StoreBufferU32:
case Opcode::BufferAtomicIAdd32:
case Opcode::BufferAtomicSMin32:

View File

@ -79,19 +79,13 @@ OPCODE(LoadBufferF32, F32, Opaq
OPCODE(LoadBufferF32x2, F32x2, Opaque, Opaque, )
OPCODE(LoadBufferF32x3, F32x3, Opaque, Opaque, )
OPCODE(LoadBufferF32x4, F32x4, Opaque, Opaque, )
OPCODE(LoadBufferFormatF32, F32, Opaque, Opaque, )
OPCODE(LoadBufferFormatF32x2, F32x2, Opaque, Opaque, )
OPCODE(LoadBufferFormatF32x3, F32x3, Opaque, Opaque, )
OPCODE(LoadBufferFormatF32x4, F32x4, Opaque, Opaque, )
OPCODE(LoadBufferFormatF32, F32x4, Opaque, Opaque, )
OPCODE(LoadBufferU32, U32, Opaque, Opaque, )
OPCODE(StoreBufferF32, Void, Opaque, Opaque, F32, )
OPCODE(StoreBufferF32x2, Void, Opaque, Opaque, F32x2, )
OPCODE(StoreBufferF32x3, Void, Opaque, Opaque, F32x3, )
OPCODE(StoreBufferF32x4, Void, Opaque, Opaque, F32x4, )
OPCODE(StoreBufferFormatF32, Void, Opaque, Opaque, F32, )
OPCODE(StoreBufferFormatF32x2, Void, Opaque, Opaque, F32x2, )
OPCODE(StoreBufferFormatF32x3, Void, Opaque, Opaque, F32x3, )
OPCODE(StoreBufferFormatF32x4, Void, Opaque, Opaque, F32x4, )
OPCODE(StoreBufferFormatF32, Void, Opaque, Opaque, F32x4, )
OPCODE(StoreBufferU32, Void, Opaque, Opaque, U32, )
// Buffer atomic operations

View File

@ -3,6 +3,7 @@
#include <algorithm>
#include <boost/container/small_vector.hpp>
#include "common/alignment.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/ir/breadth_first_search.h"
#include "shader_recompiler/ir/ir_emitter.h"
@ -45,10 +46,6 @@ bool IsBufferStore(const IR::Inst& inst) {
case IR::Opcode::StoreBufferF32x2:
case IR::Opcode::StoreBufferF32x3:
case IR::Opcode::StoreBufferF32x4:
case IR::Opcode::StoreBufferFormatF32:
case IR::Opcode::StoreBufferFormatF32x2:
case IR::Opcode::StoreBufferFormatF32x3:
case IR::Opcode::StoreBufferFormatF32x4:
case IR::Opcode::StoreBufferU32:
return true;
default:
@ -62,10 +59,6 @@ bool IsBufferInstruction(const IR::Inst& inst) {
case IR::Opcode::LoadBufferF32x2:
case IR::Opcode::LoadBufferF32x3:
case IR::Opcode::LoadBufferF32x4:
case IR::Opcode::LoadBufferFormatF32:
case IR::Opcode::LoadBufferFormatF32x2:
case IR::Opcode::LoadBufferFormatF32x3:
case IR::Opcode::LoadBufferFormatF32x4:
case IR::Opcode::LoadBufferU32:
case IR::Opcode::ReadConstBuffer:
case IR::Opcode::ReadConstBufferU32:
@ -75,6 +68,11 @@ bool IsBufferInstruction(const IR::Inst& inst) {
}
}
bool IsTextureBufferInstruction(const IR::Inst& inst) {
return inst.GetOpcode() == IR::Opcode::LoadBufferFormatF32 ||
inst.GetOpcode() == IR::Opcode::StoreBufferFormatF32;
}
static bool UseFP16(AmdGpu::DataFormat data_format, AmdGpu::NumberFormat num_format) {
switch (num_format) {
case AmdGpu::NumberFormat::Float:
@ -100,28 +98,6 @@ static bool UseFP16(AmdGpu::DataFormat data_format, AmdGpu::NumberFormat num_for
IR::Type BufferDataType(const IR::Inst& inst, AmdGpu::NumberFormat num_format) {
switch (inst.GetOpcode()) {
case IR::Opcode::LoadBufferFormatF32:
case IR::Opcode::LoadBufferFormatF32x2:
case IR::Opcode::LoadBufferFormatF32x3:
case IR::Opcode::LoadBufferFormatF32x4:
case IR::Opcode::StoreBufferFormatF32:
case IR::Opcode::StoreBufferFormatF32x2:
case IR::Opcode::StoreBufferFormatF32x3:
case IR::Opcode::StoreBufferFormatF32x4:
switch (num_format) {
case AmdGpu::NumberFormat::Unorm:
case AmdGpu::NumberFormat::Snorm:
case AmdGpu::NumberFormat::Uscaled:
case AmdGpu::NumberFormat::Sscaled:
case AmdGpu::NumberFormat::Uint:
case AmdGpu::NumberFormat::Sint:
case AmdGpu::NumberFormat::SnormNz:
return IR::Type::U32;
case AmdGpu::NumberFormat::Float:
return IR::Type::F32;
default:
UNREACHABLE();
}
case IR::Opcode::LoadBufferF32:
case IR::Opcode::LoadBufferF32x2:
case IR::Opcode::LoadBufferF32x3:
@ -143,20 +119,8 @@ IR::Type BufferDataType(const IR::Inst& inst, AmdGpu::NumberFormat num_format) {
}
}
bool IsImageInstruction(const IR::Inst& inst) {
bool IsImageAtomicInstruction(const IR::Inst& inst) {
switch (inst.GetOpcode()) {
case IR::Opcode::ImageSampleExplicitLod:
case IR::Opcode::ImageSampleImplicitLod:
case IR::Opcode::ImageSampleDrefExplicitLod:
case IR::Opcode::ImageSampleDrefImplicitLod:
case IR::Opcode::ImageFetch:
case IR::Opcode::ImageGather:
case IR::Opcode::ImageGatherDref:
case IR::Opcode::ImageQueryDimensions:
case IR::Opcode::ImageQueryLod:
case IR::Opcode::ImageGradient:
case IR::Opcode::ImageRead:
case IR::Opcode::ImageWrite:
case IR::Opcode::ImageAtomicIAdd32:
case IR::Opcode::ImageAtomicSMin32:
case IR::Opcode::ImageAtomicUMin32:
@ -178,20 +142,27 @@ bool IsImageStorageInstruction(const IR::Inst& inst) {
switch (inst.GetOpcode()) {
case IR::Opcode::ImageWrite:
case IR::Opcode::ImageRead:
case IR::Opcode::ImageAtomicIAdd32:
case IR::Opcode::ImageAtomicSMin32:
case IR::Opcode::ImageAtomicUMin32:
case IR::Opcode::ImageAtomicSMax32:
case IR::Opcode::ImageAtomicUMax32:
case IR::Opcode::ImageAtomicInc32:
case IR::Opcode::ImageAtomicDec32:
case IR::Opcode::ImageAtomicAnd32:
case IR::Opcode::ImageAtomicOr32:
case IR::Opcode::ImageAtomicXor32:
case IR::Opcode::ImageAtomicExchange32:
return true;
default:
return false;
return IsImageAtomicInstruction(inst);
}
}
bool IsImageInstruction(const IR::Inst& inst) {
switch (inst.GetOpcode()) {
case IR::Opcode::ImageSampleExplicitLod:
case IR::Opcode::ImageSampleImplicitLod:
case IR::Opcode::ImageSampleDrefExplicitLod:
case IR::Opcode::ImageSampleDrefImplicitLod:
case IR::Opcode::ImageFetch:
case IR::Opcode::ImageGather:
case IR::Opcode::ImageGatherDref:
case IR::Opcode::ImageQueryDimensions:
case IR::Opcode::ImageQueryLod:
case IR::Opcode::ImageGradient:
return true;
default:
return IsImageStorageInstruction(inst);
}
}
@ -214,7 +185,8 @@ u32 ImageOffsetArgumentPosition(const IR::Inst& inst) {
class Descriptors {
public:
explicit Descriptors(Info& info_)
: info{info_}, buffer_resources{info_.buffers}, image_resources{info_.images},
: info{info_}, buffer_resources{info_.buffers},
texture_buffer_resources{info_.texture_buffers}, image_resources{info_.images},
sampler_resources{info_.samplers} {}
u32 Add(const BufferResource& desc) {
@ -224,13 +196,21 @@ public:
desc.inline_cbuf == existing.inline_cbuf;
})};
auto& buffer = buffer_resources[index];
ASSERT(buffer.length == desc.length);
buffer.is_storage |= desc.is_storage;
buffer.used_types |= desc.used_types;
buffer.is_written |= desc.is_written;
return index;
}
u32 Add(const TextureBufferResource& desc) {
const u32 index{Add(texture_buffer_resources, desc, [&desc](const auto& existing) {
return desc.sgpr_base == existing.sgpr_base &&
desc.dword_offset == existing.dword_offset;
})};
auto& buffer = texture_buffer_resources[index];
buffer.is_written |= desc.is_written;
return index;
}
u32 Add(const ImageResource& desc) {
const u32 index{Add(image_resources, desc, [&desc](const auto& existing) {
return desc.sgpr_base == existing.sgpr_base &&
@ -247,7 +227,7 @@ public:
return true;
}
// Samplers with different bindings might still be the same.
return existing.GetSsharp(info) == desc.GetSsharp(info);
return existing.GetSharp(info) == desc.GetSharp(info);
})};
return index;
}
@ -265,6 +245,7 @@ private:
const Info& info;
BufferResourceList& buffer_resources;
TextureBufferResourceList& texture_buffer_resources;
ImageResourceList& image_resources;
SamplerResourceList& sampler_resources;
};
@ -361,33 +342,6 @@ SharpLocation TrackSharp(const IR::Inst* inst) {
};
}
static constexpr size_t MaxUboSize = 65536;
static bool IsLoadBufferFormat(const IR::Inst& inst) {
switch (inst.GetOpcode()) {
case IR::Opcode::LoadBufferFormatF32:
case IR::Opcode::LoadBufferFormatF32x2:
case IR::Opcode::LoadBufferFormatF32x3:
case IR::Opcode::LoadBufferFormatF32x4:
return true;
default:
return false;
}
}
static u32 BufferLength(const AmdGpu::Buffer& buffer) {
const auto stride = buffer.GetStride();
if (stride < sizeof(f32)) {
ASSERT(sizeof(f32) % stride == 0);
return (((buffer.num_records - 1) / sizeof(f32)) + 1) * stride;
} else if (stride == sizeof(f32)) {
return buffer.num_records;
} else {
ASSERT(stride % sizeof(f32) == 0);
return buffer.num_records * (stride / sizeof(f32));
}
}
s32 TryHandleInlineCbuf(IR::Inst& inst, Info& info, Descriptors& descriptors,
AmdGpu::Buffer& cbuf) {
@ -414,10 +368,8 @@ s32 TryHandleInlineCbuf(IR::Inst& inst, Info& info, Descriptors& descriptors,
return descriptors.Add(BufferResource{
.sgpr_base = std::numeric_limits<u32>::max(),
.dword_offset = 0,
.length = BufferLength(cbuf),
.used_types = BufferDataType(inst, cbuf.GetNumberFmt()),
.inline_cbuf = cbuf,
.is_storage = IsBufferStore(inst) || cbuf.GetSize() > MaxUboSize,
});
}
@ -429,28 +381,17 @@ void PatchBufferInstruction(IR::Block& block, IR::Inst& inst, Info& info,
IR::Inst* handle = inst.Arg(0).InstRecursive();
IR::Inst* producer = handle->Arg(0).InstRecursive();
const auto sharp = TrackSharp(producer);
const bool is_store = IsBufferStore(inst);
buffer = info.ReadUd<AmdGpu::Buffer>(sharp.sgpr_base, sharp.dword_offset);
binding = descriptors.Add(BufferResource{
.sgpr_base = sharp.sgpr_base,
.dword_offset = sharp.dword_offset,
.length = BufferLength(buffer),
.used_types = BufferDataType(inst, buffer.GetNumberFmt()),
.is_storage = is_store || buffer.GetSize() > MaxUboSize,
.is_written = is_store,
.is_written = IsBufferStore(inst),
});
}
// Update buffer descriptor format.
const auto inst_info = inst.Flags<IR::BufferInstInfo>();
auto& buffer_desc = info.buffers[binding];
if (inst_info.is_typed) {
buffer_desc.dfmt = inst_info.dmft;
buffer_desc.nfmt = inst_info.nfmt;
} else {
buffer_desc.dfmt = buffer.GetDataFmt();
buffer_desc.nfmt = buffer.GetNumberFmt();
}
// Replace handle with binding index in buffer resource list.
IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
@ -463,20 +404,7 @@ void PatchBufferInstruction(IR::Block& block, IR::Inst& inst, Info& info,
return;
}
if (IsLoadBufferFormat(inst)) {
if (UseFP16(buffer.GetDataFmt(), buffer.GetNumberFmt())) {
info.uses_fp16 = true;
}
} else {
const u32 stride = buffer.GetStride();
if (stride < 4) {
LOG_WARNING(Render_Vulkan,
"non-formatting load_buffer_* is not implemented for stride {}", stride);
}
}
// Compute address of the buffer using the stride.
// Todo: What if buffer is rebound with different stride?
IR::U32 address = ir.Imm32(inst_info.inst_offset.Value());
if (inst_info.index_enable) {
const IR::U32 index = inst_info.offset_enable ? IR::U32{ir.CompositeExtract(inst.Arg(1), 0)}
@ -491,8 +419,31 @@ void PatchBufferInstruction(IR::Block& block, IR::Inst& inst, Info& info,
inst.SetArg(1, address);
}
void PatchTextureBufferInstruction(IR::Block& block, IR::Inst& inst, Info& info,
Descriptors& descriptors) {
const IR::Inst* handle = inst.Arg(0).InstRecursive();
const IR::Inst* producer = handle->Arg(0).InstRecursive();
const auto sharp = TrackSharp(producer);
const auto buffer = info.ReadUd<AmdGpu::Buffer>(sharp.sgpr_base, sharp.dword_offset);
const s32 binding = descriptors.Add(TextureBufferResource{
.sgpr_base = sharp.sgpr_base,
.dword_offset = sharp.dword_offset,
.nfmt = buffer.GetNumberFmt(),
.is_written = inst.GetOpcode() == IR::Opcode::StoreBufferFormatF32,
});
// Replace handle with binding index in texture buffer resource list.
IR::IREmitter ir{block, IR::Block::InstructionList::s_iterator_to(inst)};
inst.SetArg(0, ir.Imm32(binding));
ASSERT(!buffer.swizzle_enable && !buffer.add_tid_enable);
}
IR::Value PatchCubeCoord(IR::IREmitter& ir, const IR::Value& s, const IR::Value& t,
const IR::Value& z) {
const IR::Value& z, bool is_storage) {
// When cubemap is written with imageStore it is treated like 2DArray.
if (is_storage) {
return ir.CompositeConstruct(s, t, z);
}
// We need to fix x and y coordinate,
// because the s and t coordinate will be scaled and plus 1.5 by v_madak_f32.
// We already force the scale value to be 1.0 when handling v_cubema_f32,
@ -530,13 +481,15 @@ void PatchImageInstruction(IR::Block& block, IR::Inst& inst, Info& info, Descrip
return;
}
ASSERT(image.GetType() != AmdGpu::ImageType::Invalid);
const bool is_storage = IsImageStorageInstruction(inst);
u32 image_binding = descriptors.Add(ImageResource{
.sgpr_base = tsharp.sgpr_base,
.dword_offset = tsharp.dword_offset,
.type = image.GetType(),
.nfmt = static_cast<AmdGpu::NumberFormat>(image.GetNumberFmt()),
.is_storage = IsImageStorageInstruction(inst),
.is_storage = is_storage,
.is_depth = bool(inst_info.is_depth),
.is_atomic = IsImageAtomicInstruction(inst),
});
// Read sampler sharp. This doesn't exist for IMAGE_LOAD/IMAGE_STORE instructions
@ -593,7 +546,8 @@ void PatchImageInstruction(IR::Block& block, IR::Inst& inst, Info& info, Descrip
case AmdGpu::ImageType::Color3D: // x, y, z
return {ir.CompositeConstruct(body->Arg(0), body->Arg(1), body->Arg(2)), body->Arg(3)};
case AmdGpu::ImageType::Cube: // x, y, face
return {PatchCubeCoord(ir, body->Arg(0), body->Arg(1), body->Arg(2)), body->Arg(3)};
return {PatchCubeCoord(ir, body->Arg(0), body->Arg(1), body->Arg(2), is_storage),
body->Arg(3)};
default:
UNREACHABLE_MSG("Unknown image type {}", image.GetType());
}
@ -668,6 +622,10 @@ void ResourceTrackingPass(IR::Program& program) {
PatchBufferInstruction(*block, inst, info, descriptors);
continue;
}
if (IsTextureBufferInstruction(inst)) {
PatchTextureBufferInstruction(*block, inst, info, descriptors);
continue;
}
if (IsImageInstruction(inst)) {
PatchImageInstruction(*block, inst, info, descriptors);
}

View File

@ -29,6 +29,12 @@ void Visit(Info& info, IR::Inst& inst) {
case IR::Opcode::ImageWrite:
info.has_storage_images = true;
break;
case IR::Opcode::LoadBufferFormatF32:
info.has_texel_buffers = true;
break;
case IR::Opcode::StoreBufferFormatF32:
info.has_image_buffers = true;
break;
case IR::Opcode::QuadShuffle:
info.uses_group_quad = true;
break;
@ -44,6 +50,9 @@ void Visit(Info& info, IR::Inst& inst) {
case IR::Opcode::ImageQueryLod:
info.has_image_query = true;
break;
case IR::Opcode::LaneId:
info.uses_lane_id = true;
break;
default:
break;
}

View File

@ -12,11 +12,13 @@
namespace Shader::IR {
struct Program {
explicit Program(Info& info_) : info{info_} {}
AbstractSyntaxList syntax_list;
BlockList blocks;
BlockList post_order_blocks;
std::vector<Gcn::GcnInst> ins_list;
Info info;
Info& info;
};
[[nodiscard]] std::string DumpProgram(const Program& program);

View File

@ -66,9 +66,6 @@ union BufferInstInfo {
BitField<0, 1, u32> index_enable;
BitField<1, 1, u32> offset_enable;
BitField<2, 12, u32> inst_offset;
BitField<14, 4, AmdGpu::DataFormat> dmft;
BitField<18, 3, AmdGpu::NumberFormat> nfmt;
BitField<21, 1, u32> is_typed;
};
enum class ScalarReg : u32 {

View File

@ -29,7 +29,7 @@ IR::BlockList GenerateBlocks(const IR::AbstractSyntaxList& syntax_list) {
IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool, std::span<const u32> token,
const Info&& info, const Profile& profile) {
Info& info, const Profile& profile) {
// Ensure first instruction is expected.
constexpr u32 token_mov_vcchi = 0xBEEB03FF;
ASSERT_MSG(token[0] == token_mov_vcchi, "First instruction is not s_mov_b32 vcc_hi, #imm");
@ -38,7 +38,7 @@ IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Gcn::GcnDecodeContext decoder;
// Decode and save instructions
IR::Program program;
IR::Program program{info};
program.ins_list.reserve(token.size());
while (!slice.atEnd()) {
program.ins_list.emplace_back(decoder.decodeInstruction(slice));
@ -49,7 +49,6 @@ IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Gcn::CFG cfg{gcn_block_pool, program.ins_list};
// Structurize control flow graph and create program.
program.info = std::move(info);
program.syntax_list = Shader::Gcn::BuildASL(inst_pool, block_pool, cfg, program.info, profile);
program.blocks = GenerateBlocks(program.syntax_list);
program.post_order_blocks = Shader::IR::PostOrder(program.syntax_list.front());

View File

@ -13,7 +13,7 @@ struct Profile;
[[nodiscard]] IR::Program TranslateProgram(Common::ObjectPool<IR::Inst>& inst_pool,
Common::ObjectPool<IR::Block>& block_pool,
std::span<const u32> code, const Info&& info,
std::span<const u32> code, Info& info,
const Profile& profile);
} // namespace Shader

View File

@ -4,6 +4,7 @@
#pragma once
#include <span>
#include <boost/container/small_vector.hpp>
#include <boost/container/static_vector.hpp>
#include "common/assert.h"
#include "common/types.h"
@ -74,18 +75,29 @@ struct Info;
struct BufferResource {
u32 sgpr_base;
u32 dword_offset;
u32 length;
IR::Type used_types;
AmdGpu::Buffer inline_cbuf;
AmdGpu::DataFormat dfmt;
AmdGpu::NumberFormat nfmt;
bool is_storage{};
bool is_instance_data{};
bool is_written{};
constexpr AmdGpu::Buffer GetVsharp(const Info& info) const noexcept;
bool IsStorage(AmdGpu::Buffer buffer) const noexcept {
static constexpr size_t MaxUboSize = 65536;
return buffer.GetSize() > MaxUboSize || is_written;
}
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using BufferResourceList = boost::container::static_vector<BufferResource, 16>;
using BufferResourceList = boost::container::small_vector<BufferResource, 16>;
struct TextureBufferResource {
u32 sgpr_base;
u32 dword_offset;
AmdGpu::NumberFormat nfmt;
bool is_written{};
constexpr AmdGpu::Buffer GetSharp(const Info& info) const noexcept;
};
using TextureBufferResourceList = boost::container::small_vector<TextureBufferResource, 16>;
struct ImageResource {
u32 sgpr_base;
@ -94,8 +106,11 @@ struct ImageResource {
AmdGpu::NumberFormat nfmt;
bool is_storage;
bool is_depth;
bool is_atomic{};
constexpr AmdGpu::Image GetSharp(const Info& info) const noexcept;
};
using ImageResourceList = boost::container::static_vector<ImageResource, 16>;
using ImageResourceList = boost::container::small_vector<ImageResource, 16>;
struct SamplerResource {
u32 sgpr_base;
@ -104,9 +119,9 @@ struct SamplerResource {
u32 associated_image : 4;
u32 disable_aniso : 1;
constexpr AmdGpu::Sampler GetSsharp(const Info& info) const noexcept;
constexpr AmdGpu::Sampler GetSharp(const Info& info) const noexcept;
};
using SamplerResourceList = boost::container::static_vector<SamplerResource, 16>;
using SamplerResourceList = boost::container::small_vector<SamplerResource, 16>;
struct PushData {
static constexpr size_t BufOffsetIndex = 2;
@ -179,6 +194,7 @@ struct Info {
s8 instance_offset_sgpr = -1;
BufferResourceList buffers;
TextureBufferResourceList texture_buffers;
ImageResourceList images;
SamplerResourceList samplers;
@ -194,9 +210,12 @@ struct Info {
u64 pgm_hash{};
u32 shared_memory_size{};
bool has_storage_images{};
bool has_image_buffers{};
bool has_texel_buffers{};
bool has_discard{};
bool has_image_gather{};
bool has_image_query{};
bool uses_lane_id{};
bool uses_group_quad{};
bool uses_shared{};
bool uses_fp16{};
@ -214,6 +233,10 @@ struct Info {
return data;
}
size_t NumBindings() const noexcept {
return buffers.size() + texture_buffers.size() + images.size() + samplers.size();
}
[[nodiscard]] std::pair<u32, u32> GetDrawOffsets() const noexcept {
u32 vertex_offset = 0;
u32 instance_offset = 0;
@ -227,11 +250,19 @@ struct Info {
}
};
constexpr AmdGpu::Buffer BufferResource::GetVsharp(const Info& info) const noexcept {
constexpr AmdGpu::Buffer BufferResource::GetSharp(const Info& info) const noexcept {
return inline_cbuf ? inline_cbuf : info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Sampler SamplerResource::GetSsharp(const Info& info) const noexcept {
constexpr AmdGpu::Buffer TextureBufferResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Buffer>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Image ImageResource::GetSharp(const Info& info) const noexcept {
return info.ReadUd<AmdGpu::Image>(sgpr_base, dword_offset);
}
constexpr AmdGpu::Sampler SamplerResource::GetSharp(const Info& info) const noexcept {
return inline_sampler ? inline_sampler : info.ReadUd<AmdGpu::Sampler>(sgpr_base, dword_offset);
}

View File

@ -383,6 +383,22 @@ Liverpool::Task Liverpool::ProcessGraphics(std::span<const u32> dcb, std::span<c
}
break;
}
case PM4ItOpcode::DispatchIndirect: {
const auto* dispatch_indirect =
reinterpret_cast<const PM4CmdDispatchIndirect*>(header);
const auto offset = dispatch_indirect->data_offset;
const auto ib_address = mapped_queues[GfxQueueId].indirect_args_addr;
const auto size = sizeof(PM4CmdDispatchIndirect::GroupDimensions);
if (rasterizer && (regs.cs_program.dispatch_initiator & 1)) {
const auto cmd_address = reinterpret_cast<const void*>(header);
rasterizer->ScopeMarkerBegin(
fmt::format("dcb:{}:DispatchIndirect", cmd_address));
rasterizer->Breadcrumb(u64(cmd_address));
rasterizer->DispatchIndirect(ib_address, offset, size);
rasterizer->ScopeMarkerEnd();
}
break;
}
case PM4ItOpcode::NumInstances: {
const auto* num_instances = reinterpret_cast<const PM4CmdDrawNumInstances*>(header);
regs.num_instances.num_instances = num_instances->num_instances;
@ -399,6 +415,12 @@ Liverpool::Task Liverpool::ProcessGraphics(std::span<const u32> dcb, std::span<c
regs.num_indices = index_size->num_indices;
break;
}
case PM4ItOpcode::SetBase: {
const auto* set_base = reinterpret_cast<const PM4CmdSetBase*>(header);
ASSERT(set_base->base_index == PM4CmdSetBase::BaseIndex::DrawIndexIndirPatchTable);
mapped_queues[GfxQueueId].indirect_args_addr = set_base->Address<u64>();
break;
}
case PM4ItOpcode::EventWrite: {
// const auto* event = reinterpret_cast<const PM4CmdEventWrite*>(header);
break;

View File

@ -167,7 +167,7 @@ struct Liverpool {
static constexpr auto* GetBinaryInfo(const Shader& sh) {
const auto* code = sh.template Address<u32*>();
const auto* bininfo = std::bit_cast<const BinaryInfo*>(code + (code[1] + 1) * 2);
ASSERT_MSG(bininfo->Valid(), "Invalid shader binary header");
// ASSERT_MSG(bininfo->Valid(), "Invalid shader binary header");
return bininfo;
}
@ -1127,6 +1127,7 @@ private:
std::vector<u32> ccb_buffer;
std::queue<Task::Handle> submits{};
ComputeProgram cs_state{};
VAddr indirect_args_addr{};
};
std::array<GpuQueue, NumTotalQueues> mapped_queues{};

View File

@ -61,6 +61,10 @@ enum class NumberFormat : u32 {
Ubscaled = 13,
};
[[nodiscard]] constexpr bool IsInteger(NumberFormat nfmt) {
return nfmt == AmdGpu::NumberFormat::Sint || nfmt == AmdGpu::NumberFormat::Uint;
}
[[nodiscard]] std::string_view NameOf(DataFormat fmt);
[[nodiscard]] std::string_view NameOf(NumberFormat fmt);

View File

@ -704,4 +704,40 @@ struct PM4CmdReleaseMem {
}
};
struct PM4CmdSetBase {
enum class BaseIndex : u32 {
DisplayListPatchTable = 0b0000,
DrawIndexIndirPatchTable = 0b0001,
GdsPartition = 0b0010,
CePartition = 0b0011,
};
PM4Type3Header header;
union {
BitField<0, 4, BaseIndex> base_index;
u32 dw1;
};
u32 address0;
u32 address1;
template <typename T>
T Address() const {
ASSERT(base_index == BaseIndex::DisplayListPatchTable ||
base_index == BaseIndex::DrawIndexIndirPatchTable);
return reinterpret_cast<T>(address0 | (u64(address1 & 0xffff) << 32u));
}
};
struct PM4CmdDispatchIndirect {
struct GroupDimensions {
u32 dim_x;
u32 dim_y;
u32 dim_z;
};
PM4Type3Header header;
u32 data_offset; ///< Byte aligned offset where the required data structure starts
u32 dispatch_initiator; ///< Dispatch Initiator Register
};
} // namespace AmdGpu

View File

@ -3,6 +3,7 @@
#pragma once
#include "common/alignment.h"
#include "common/assert.h"
#include "common/bit_field.h"
#include "common/types.h"
@ -68,6 +69,10 @@ struct Buffer {
return stride == 0 ? 1U : stride;
}
u32 NumDwords() const noexcept {
return Common::AlignUp(GetSize(), sizeof(u32)) >> 2;
}
u32 GetSize() const noexcept {
return GetStride() * num_records;
}

View File

@ -13,12 +13,6 @@
namespace VideoCore {
constexpr vk::BufferUsageFlags AllFlags =
vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eTransferDst |
vk::BufferUsageFlagBits::eUniformTexelBuffer | vk::BufferUsageFlagBits::eStorageTexelBuffer |
vk::BufferUsageFlagBits::eUniformBuffer | vk::BufferUsageFlagBits::eStorageBuffer |
vk::BufferUsageFlagBits::eIndexBuffer | vk::BufferUsageFlagBits::eVertexBuffer;
std::string_view BufferTypeName(MemoryUsage type) {
switch (type) {
case MemoryUsage::Upload:
@ -95,13 +89,13 @@ void UniqueBuffer::Create(const vk::BufferCreateInfo& buffer_ci, MemoryUsage usa
}
Buffer::Buffer(const Vulkan::Instance& instance_, MemoryUsage usage_, VAddr cpu_addr_,
u64 size_bytes_)
vk::BufferUsageFlags flags, u64 size_bytes_)
: cpu_addr{cpu_addr_}, size_bytes{size_bytes_}, instance{&instance_}, usage{usage_},
buffer{instance->GetDevice(), instance->GetAllocator()} {
// Create buffer object.
const vk::BufferCreateInfo buffer_ci = {
.size = size_bytes,
.usage = AllFlags,
.usage = flags,
};
VmaAllocationInfo alloc_info{};
buffer.Create(buffer_ci, usage, &alloc_info);
@ -118,27 +112,35 @@ Buffer::Buffer(const Vulkan::Instance& instance_, MemoryUsage usage_, VAddr cpu_
is_coherent = property_flags & VK_MEMORY_PROPERTY_HOST_COHERENT_BIT;
}
vk::BufferView Buffer::View(u32 offset, u32 size, AmdGpu::DataFormat dfmt,
vk::BufferView Buffer::View(u32 offset, u32 size, bool is_written, AmdGpu::DataFormat dfmt,
AmdGpu::NumberFormat nfmt) {
const auto it{std::ranges::find_if(views, [offset, size, dfmt, nfmt](const BufferView& view) {
return offset == view.offset && size == view.size && dfmt == view.dfmt && nfmt == view.nfmt;
const auto it{std::ranges::find_if(views, [=](const BufferView& view) {
return offset == view.offset && size == view.size && is_written == view.is_written &&
dfmt == view.dfmt && nfmt == view.nfmt;
})};
if (it != views.end()) {
return it->handle;
return *it->handle;
}
views.push_back({
.offset = offset,
.size = size,
.dfmt = dfmt,
.nfmt = nfmt,
.handle = instance->GetDevice().createBufferView({
const vk::BufferUsageFlags2CreateInfoKHR usage_flags = {
.usage = is_written ? vk::BufferUsageFlagBits2KHR::eStorageTexelBuffer
: vk::BufferUsageFlagBits2KHR::eUniformTexelBuffer,
};
const vk::BufferViewCreateInfo view_ci = {
.pNext = &usage_flags,
.buffer = buffer.buffer,
.format = Vulkan::LiverpoolToVK::SurfaceFormat(dfmt, nfmt),
.offset = offset,
.range = size,
}),
};
views.push_back({
.offset = offset,
.size = size,
.is_written = is_written,
.dfmt = dfmt,
.nfmt = nfmt,
.handle = instance->GetDevice().createBufferViewUnique(view_ci),
});
return views.back().handle;
return *views.back().handle;
}
constexpr u64 WATCHES_INITIAL_RESERVE = 0x4000;
@ -146,7 +148,7 @@ constexpr u64 WATCHES_RESERVE_CHUNK = 0x1000;
StreamBuffer::StreamBuffer(const Vulkan::Instance& instance, Vulkan::Scheduler& scheduler_,
MemoryUsage usage, u64 size_bytes)
: Buffer{instance, usage, 0, size_bytes}, scheduler{scheduler_} {
: Buffer{instance, usage, 0, AllFlags, size_bytes}, scheduler{scheduler_} {
ReserveWatches(current_watches, WATCHES_INITIAL_RESERVE);
ReserveWatches(previous_watches, WATCHES_INITIAL_RESERVE);
const auto device = instance.GetDevice();

View File

@ -31,6 +31,15 @@ enum class MemoryUsage {
Stream, ///< Requests device local host visible buffer, falling back host memory.
};
constexpr vk::BufferUsageFlags ReadFlags =
vk::BufferUsageFlagBits::eTransferSrc | vk::BufferUsageFlagBits::eUniformTexelBuffer |
vk::BufferUsageFlagBits::eUniformBuffer | vk::BufferUsageFlagBits::eIndexBuffer |
vk::BufferUsageFlagBits::eVertexBuffer | vk::BufferUsageFlagBits::eIndirectBuffer;
constexpr vk::BufferUsageFlags AllFlags = ReadFlags | vk::BufferUsageFlagBits::eTransferDst |
vk::BufferUsageFlagBits::eStorageTexelBuffer |
vk::BufferUsageFlagBits::eStorageBuffer;
struct UniqueBuffer {
explicit UniqueBuffer(vk::Device device, VmaAllocator allocator);
~UniqueBuffer();
@ -65,7 +74,7 @@ struct UniqueBuffer {
class Buffer {
public:
explicit Buffer(const Vulkan::Instance& instance, MemoryUsage usage, VAddr cpu_addr_,
u64 size_bytes_);
vk::BufferUsageFlags flags, u64 size_bytes_);
Buffer& operator=(const Buffer&) = delete;
Buffer(const Buffer&) = delete;
@ -73,7 +82,8 @@ public:
Buffer& operator=(Buffer&&) = default;
Buffer(Buffer&&) = default;
vk::BufferView View(u32 offset, u32 size, AmdGpu::DataFormat dfmt, AmdGpu::NumberFormat nfmt);
vk::BufferView View(u32 offset, u32 size, bool is_written, AmdGpu::DataFormat dfmt,
AmdGpu::NumberFormat nfmt);
/// Increases the likeliness of this being a stream buffer
void IncreaseStreamScore(int score) noexcept {
@ -121,9 +131,10 @@ public:
struct BufferView {
u32 offset;
u32 size;
bool is_written;
AmdGpu::DataFormat dfmt;
AmdGpu::NumberFormat nfmt;
vk::BufferView handle;
vk::UniqueBufferView handle;
};
std::vector<BufferView> views;
};

View File

@ -23,7 +23,7 @@ BufferCache::BufferCache(const Vulkan::Instance& instance_, Vulkan::Scheduler& s
stream_buffer{instance, scheduler, MemoryUsage::Stream, UboStreamBufferSize},
memory_tracker{&tracker} {
// Ensure the first slot is used for the null buffer
void(slot_buffers.insert(instance, MemoryUsage::DeviceLocal, 0, 1));
void(slot_buffers.insert(instance, MemoryUsage::DeviceLocal, 0, ReadFlags, 1));
}
BufferCache::~BufferCache() = default;
@ -228,11 +228,12 @@ u32 BufferCache::BindIndexBuffer(bool& is_indexed, u32 index_offset) {
return regs.num_indices;
}
std::pair<Buffer*, u32> BufferCache::ObtainBuffer(VAddr device_addr, u32 size, bool is_written) {
std::pair<Buffer*, u32> BufferCache::ObtainBuffer(VAddr device_addr, u32 size, bool is_written,
bool is_texel_buffer) {
std::scoped_lock lk{mutex};
static constexpr u64 StreamThreshold = CACHING_PAGESIZE;
const bool is_gpu_dirty = memory_tracker.IsRegionGpuModified(device_addr, size);
if (!is_written && size < StreamThreshold && !is_gpu_dirty) {
if (!is_written && !is_texel_buffer && size <= StreamThreshold && !is_gpu_dirty) {
// For small uniform buffers that have not been modified by gpu
// use device local stream buffer to reduce renderpass breaks.
const u64 offset = stream_buffer.Copy(device_addr, size, instance.UniformMinAlignment());
@ -421,7 +422,7 @@ BufferId BufferCache::CreateBuffer(VAddr device_addr, u32 wanted_size) {
const OverlapResult overlap = ResolveOverlaps(device_addr, wanted_size);
const u32 size = static_cast<u32>(overlap.end - overlap.begin);
const BufferId new_buffer_id =
slot_buffers.insert(instance, MemoryUsage::DeviceLocal, overlap.begin, size);
slot_buffers.insert(instance, MemoryUsage::DeviceLocal, overlap.begin, AllFlags, size);
auto& new_buffer = slot_buffers[new_buffer_id];
const size_t size_bytes = new_buffer.SizeBytes();
const auto cmdbuf = scheduler.CommandBuffer();
@ -495,7 +496,8 @@ bool BufferCache::SynchronizeBuffer(Buffer& buffer, VAddr device_addr, u32 size)
} else {
// For large one time transfers use a temporary host buffer.
// RenderDoc can lag quite a bit if the stream buffer is too large.
Buffer temp_buffer{instance, MemoryUsage::Upload, 0, total_size_bytes};
Buffer temp_buffer{instance, MemoryUsage::Upload, 0, vk::BufferUsageFlagBits::eTransferSrc,
total_size_bytes};
src_buffer = temp_buffer.Handle();
u8* const staging = temp_buffer.mapped_data.data();
for (auto& copy : copies) {

View File

@ -66,7 +66,8 @@ public:
u32 BindIndexBuffer(bool& is_indexed, u32 index_offset);
/// Obtains a buffer for the specified region.
[[nodiscard]] std::pair<Buffer*, u32> ObtainBuffer(VAddr gpu_addr, u32 size, bool is_written);
[[nodiscard]] std::pair<Buffer*, u32> ObtainBuffer(VAddr gpu_addr, u32 size, bool is_written,
bool is_texel_buffer = false);
/// Obtains a temporary buffer for usage in texture cache.
[[nodiscard]] std::pair<const Buffer*, u32> ObtainTempBuffer(VAddr gpu_addr, u32 size);

View File

@ -13,26 +13,35 @@ namespace Vulkan {
ComputePipeline::ComputePipeline(const Instance& instance_, Scheduler& scheduler_,
vk::PipelineCache pipeline_cache, u64 compute_key_,
const Program* program)
: instance{instance_}, scheduler{scheduler_}, compute_key{compute_key_},
info{&program->pgm.info} {
const Shader::Info& info_, vk::ShaderModule module)
: instance{instance_}, scheduler{scheduler_}, compute_key{compute_key_}, info{&info_} {
const vk::PipelineShaderStageCreateInfo shader_ci = {
.stage = vk::ShaderStageFlagBits::eCompute,
.module = program->module,
.module = module,
.pName = "main",
};
u32 binding{};
boost::container::small_vector<vk::DescriptorSetLayoutBinding, 32> bindings;
for (const auto& buffer : info->buffers) {
const auto sharp = buffer.GetSharp(*info);
bindings.push_back({
.binding = binding++,
.descriptorType = buffer.is_storage ? vk::DescriptorType::eStorageBuffer
.descriptorType = buffer.IsStorage(sharp) ? vk::DescriptorType::eStorageBuffer
: vk::DescriptorType::eUniformBuffer,
.descriptorCount = 1,
.stageFlags = vk::ShaderStageFlagBits::eCompute,
});
}
for (const auto& tex_buffer : info->texture_buffers) {
bindings.push_back({
.binding = binding++,
.descriptorType = tex_buffer.is_written ? vk::DescriptorType::eStorageTexelBuffer
: vk::DescriptorType::eUniformTexelBuffer,
.descriptorCount = 1,
.stageFlags = vk::ShaderStageFlagBits::eCompute,
});
}
for (const auto& image : info->images) {
bindings.push_back({
.binding = binding++,
@ -91,22 +100,24 @@ ComputePipeline::~ComputePipeline() = default;
bool ComputePipeline::BindResources(VideoCore::BufferCache& buffer_cache,
VideoCore::TextureCache& texture_cache) const {
// Bind resource buffers and textures.
boost::container::static_vector<vk::BufferView, 8> buffer_views;
boost::container::static_vector<vk::DescriptorBufferInfo, 16> buffer_infos;
boost::container::static_vector<vk::DescriptorImageInfo, 16> image_infos;
boost::container::small_vector<vk::WriteDescriptorSet, 16> set_writes;
Shader::PushData push_data{};
u32 binding{};
for (const auto& buffer : info->buffers) {
const auto vsharp = buffer.GetVsharp(*info);
for (const auto& desc : info->buffers) {
const auto vsharp = desc.GetSharp(*info);
const bool is_storage = desc.IsStorage(vsharp);
const VAddr address = vsharp.base_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 (buffer.is_storage) {
if (desc.is_written) {
if (texture_cache.TouchMeta(address, true)) {
LOG_WARNING(Render_Vulkan, "Metadata update skipped");
LOG_TRACE(Render_Vulkan, "Metadata update skipped");
return false;
}
} else {
@ -115,13 +126,12 @@ bool ComputePipeline::BindResources(VideoCore::BufferCache& buffer_cache,
}
}
const u32 size = vsharp.GetSize();
if (buffer.is_written) {
texture_cache.InvalidateMemory(address, size, true);
if (desc.is_written) {
texture_cache.InvalidateMemory(address, size);
}
const u32 alignment =
buffer.is_storage ? instance.StorageMinAlignment() : instance.UniformMinAlignment();
const auto [vk_buffer, offset] =
buffer_cache.ObtainBuffer(address, size, buffer.is_written);
is_storage ? instance.StorageMinAlignment() : instance.UniformMinAlignment();
const auto [vk_buffer, offset] = buffer_cache.ObtainBuffer(address, size, desc.is_written);
const u32 offset_aligned = Common::AlignDown(offset, alignment);
const u32 adjust = offset - offset_aligned;
if (adjust != 0) {
@ -134,20 +144,68 @@ bool ComputePipeline::BindResources(VideoCore::BufferCache& buffer_cache,
.dstBinding = binding++,
.dstArrayElement = 0,
.descriptorCount = 1,
.descriptorType = buffer.is_storage ? vk::DescriptorType::eStorageBuffer
.descriptorType = is_storage ? vk::DescriptorType::eStorageBuffer
: vk::DescriptorType::eUniformBuffer,
.pBufferInfo = &buffer_infos.back(),
});
}
for (const auto& desc : info->texture_buffers) {
const auto vsharp = desc.GetSharp(*info);
vk::BufferView& buffer_view = buffer_views.emplace_back(VK_NULL_HANDLE);
if (vsharp.GetDataFmt() != AmdGpu::DataFormat::FormatInvalid) {
const VAddr address = vsharp.base_address;
const u32 size = vsharp.GetSize();
if (desc.is_written) {
if (texture_cache.TouchMeta(address, true)) {
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 (desc.is_written) {
texture_cache.InvalidateMemory(address, size);
}
const u32 alignment = instance.TexelBufferMinAlignment();
const auto [vk_buffer, offset] =
buffer_cache.ObtainBuffer(address, size, desc.is_written, true);
const u32 fmt_stride = AmdGpu::NumBits(vsharp.GetDataFmt()) >> 3;
ASSERT_MSG(fmt_stride == vsharp.GetStride(),
"Texel buffer stride must match format stride");
const u32 offset_aligned = Common::AlignDown(offset, alignment);
const u32 adjust = offset - offset_aligned;
if (adjust != 0) {
ASSERT(adjust % fmt_stride == 0);
push_data.AddOffset(binding, adjust / fmt_stride);
}
buffer_view = vk_buffer->View(offset_aligned, size + adjust, desc.is_written,
vsharp.GetDataFmt(), vsharp.GetNumberFmt());
}
set_writes.push_back({
.dstSet = VK_NULL_HANDLE,
.dstBinding = binding++,
.dstArrayElement = 0,
.descriptorCount = 1,
.descriptorType = desc.is_written ? vk::DescriptorType::eStorageTexelBuffer
: vk::DescriptorType::eUniformTexelBuffer,
.pTexelBufferView = &buffer_view,
});
}
for (const auto& image_desc : info->images) {
const auto tsharp =
info->ReadUd<AmdGpu::Image>(image_desc.sgpr_base, image_desc.dword_offset);
const auto tsharp = image_desc.GetSharp(*info);
if (tsharp.GetDataFmt() != AmdGpu::DataFormat::FormatInvalid) {
VideoCore::ImageInfo image_info{tsharp};
VideoCore::ImageViewInfo view_info{tsharp, image_desc.is_storage};
const auto& image_view = texture_cache.FindTexture(image_info, view_info);
const auto& image = texture_cache.GetImage(image_view.image_id);
image_infos.emplace_back(VK_NULL_HANDLE, *image_view.image_view, image.layout);
} else {
image_infos.emplace_back(VK_NULL_HANDLE, VK_NULL_HANDLE, vk::ImageLayout::eGeneral);
}
set_writes.push_back({
.dstSet = VK_NULL_HANDLE,
.dstBinding = binding++,
@ -163,7 +221,7 @@ bool ComputePipeline::BindResources(VideoCore::BufferCache& buffer_cache,
}
}
for (const auto& sampler : info->samplers) {
const auto ssharp = sampler.GetSsharp(*info);
const auto ssharp = sampler.GetSharp(*info);
const auto vk_sampler = texture_cache.GetSampler(ssharp);
image_infos.emplace_back(vk_sampler, VK_NULL_HANDLE, vk::ImageLayout::eGeneral);
set_writes.push_back({

View File

@ -3,7 +3,7 @@
#pragma once
#include "shader_recompiler/ir/program.h"
#include <boost/container/small_vector.hpp>
#include "shader_recompiler/runtime_info.h"
#include "video_core/renderer_vulkan/vk_common.h"
@ -17,18 +17,11 @@ namespace Vulkan {
class Instance;
class Scheduler;
struct Program {
Shader::IR::Program pgm;
std::vector<u32> spv;
vk::ShaderModule module;
u32 end_binding;
};
class ComputePipeline {
public:
explicit ComputePipeline(const Instance& instance, Scheduler& scheduler,
vk::PipelineCache pipeline_cache, u64 compute_key,
const Program* program);
const Shader::Info& info, vk::ShaderModule module);
~ComputePipeline();
[[nodiscard]] vk::Pipeline Handle() const noexcept {

View File

@ -19,15 +19,11 @@ namespace Vulkan {
GraphicsPipeline::GraphicsPipeline(const Instance& instance_, Scheduler& scheduler_,
const GraphicsPipelineKey& key_,
vk::PipelineCache pipeline_cache,
std::span<const Program*, MaxShaderStages> programs)
std::span<const Shader::Info*, MaxShaderStages> infos,
std::span<const vk::ShaderModule> modules)
: instance{instance_}, scheduler{scheduler_}, key{key_} {
const vk::Device device = instance.GetDevice();
for (u32 i = 0; i < MaxShaderStages; i++) {
if (!programs[i]) {
continue;
}
stages[i] = &programs[i]->pgm.info;
}
std::ranges::copy(infos, stages.begin());
BuildDescSetLayout();
const vk::PushConstantRange push_constants = {
@ -194,16 +190,18 @@ GraphicsPipeline::GraphicsPipeline(const Instance& instance_, Scheduler& schedul
auto stage = u32(Shader::Stage::Vertex);
boost::container::static_vector<vk::PipelineShaderStageCreateInfo, MaxShaderStages>
shader_stages;
if (infos[stage]) {
shader_stages.emplace_back(vk::PipelineShaderStageCreateInfo{
.stage = vk::ShaderStageFlagBits::eVertex,
.module = programs[stage]->module,
.module = modules[stage],
.pName = "main",
});
}
stage = u32(Shader::Stage::Fragment);
if (programs[stage]) {
if (infos[stage]) {
shader_stages.emplace_back(vk::PipelineShaderStageCreateInfo{
.stage = vk::ShaderStageFlagBits::eFragment,
.module = programs[stage]->module,
.module = modules[stage],
.pName = "main",
});
}
@ -309,14 +307,24 @@ void GraphicsPipeline::BuildDescSetLayout() {
continue;
}
for (const auto& buffer : stage->buffers) {
const auto sharp = buffer.GetSharp(*stage);
bindings.push_back({
.binding = binding++,
.descriptorType = buffer.is_storage ? vk::DescriptorType::eStorageBuffer
.descriptorType = buffer.IsStorage(sharp) ? vk::DescriptorType::eStorageBuffer
: vk::DescriptorType::eUniformBuffer,
.descriptorCount = 1,
.stageFlags = vk::ShaderStageFlagBits::eVertex | vk::ShaderStageFlagBits::eFragment,
});
}
for (const auto& tex_buffer : stage->texture_buffers) {
bindings.push_back({
.binding = binding++,
.descriptorType = tex_buffer.is_written ? vk::DescriptorType::eStorageTexelBuffer
: vk::DescriptorType::eUniformTexelBuffer,
.descriptorCount = 1,
.stageFlags = vk::ShaderStageFlagBits::eVertex | vk::ShaderStageFlagBits::eFragment,
});
}
for (const auto& image : stage->images) {
bindings.push_back({
.binding = binding++,
@ -347,7 +355,8 @@ void GraphicsPipeline::BindResources(const Liverpool::Regs& regs,
VideoCore::BufferCache& buffer_cache,
VideoCore::TextureCache& texture_cache) const {
// Bind resource buffers and textures.
boost::container::static_vector<vk::DescriptorBufferInfo, 16> buffer_infos;
boost::container::static_vector<vk::BufferView, 8> buffer_views;
boost::container::static_vector<vk::DescriptorBufferInfo, 32> buffer_infos;
boost::container::static_vector<vk::DescriptorImageInfo, 32> image_infos;
boost::container::small_vector<vk::WriteDescriptorSet, 16> set_writes;
Shader::PushData push_data{};
@ -362,15 +371,16 @@ void GraphicsPipeline::BindResources(const Liverpool::Regs& regs,
push_data.step1 = regs.vgt_instance_step_rate_1;
}
for (const auto& buffer : stage->buffers) {
const auto vsharp = buffer.GetVsharp(*stage);
const auto vsharp = buffer.GetSharp(*stage);
const bool is_storage = buffer.IsStorage(vsharp);
if (vsharp) {
const VAddr address = vsharp.base_address;
if (texture_cache.IsMeta(address)) {
LOG_WARNING(Render_Vulkan, "Unexpected metadata read by a PS shader (buffer)");
}
const u32 size = vsharp.GetSize();
const u32 alignment = buffer.is_storage ? instance.StorageMinAlignment()
: instance.UniformMinAlignment();
const u32 alignment =
is_storage ? instance.StorageMinAlignment() : instance.UniformMinAlignment();
const auto [vk_buffer, offset] =
buffer_cache.ObtainBuffer(address, size, buffer.is_written);
const u32 offset_aligned = Common::AlignDown(offset, alignment);
@ -388,16 +398,47 @@ void GraphicsPipeline::BindResources(const Liverpool::Regs& regs,
.dstBinding = binding++,
.dstArrayElement = 0,
.descriptorCount = 1,
.descriptorType = buffer.is_storage ? vk::DescriptorType::eStorageBuffer
.descriptorType = is_storage ? vk::DescriptorType::eStorageBuffer
: vk::DescriptorType::eUniformBuffer,
.pBufferInfo = &buffer_infos.back(),
});
}
for (const auto& tex_buffer : stage->texture_buffers) {
const auto vsharp = tex_buffer.GetSharp(*stage);
vk::BufferView& buffer_view = buffer_views.emplace_back(VK_NULL_HANDLE);
if (vsharp.GetDataFmt() != AmdGpu::DataFormat::FormatInvalid) {
const VAddr address = vsharp.base_address;
const u32 size = vsharp.GetSize();
const u32 alignment = instance.TexelBufferMinAlignment();
const auto [vk_buffer, offset] =
buffer_cache.ObtainBuffer(address, size, tex_buffer.is_written, true);
const u32 fmt_stride = AmdGpu::NumBits(vsharp.GetDataFmt()) >> 3;
ASSERT_MSG(fmt_stride == vsharp.GetStride(),
"Texel buffer stride must match format stride");
const u32 offset_aligned = Common::AlignDown(offset, alignment);
const u32 adjust = offset - offset_aligned;
if (adjust != 0) {
ASSERT(adjust % fmt_stride == 0);
push_data.AddOffset(binding, adjust / fmt_stride);
}
buffer_view = vk_buffer->View(offset_aligned, size + adjust, tex_buffer.is_written,
vsharp.GetDataFmt(), vsharp.GetNumberFmt());
}
set_writes.push_back({
.dstSet = VK_NULL_HANDLE,
.dstBinding = binding++,
.dstArrayElement = 0,
.descriptorCount = 1,
.descriptorType = tex_buffer.is_written ? vk::DescriptorType::eStorageTexelBuffer
: vk::DescriptorType::eUniformTexelBuffer,
.pTexelBufferView = &buffer_view,
});
}
boost::container::static_vector<AmdGpu::Image, 16> tsharps;
for (const auto& image_desc : stage->images) {
const auto tsharp =
stage->ReadUd<AmdGpu::Image>(image_desc.sgpr_base, image_desc.dword_offset);
const auto tsharp = image_desc.GetSharp(*stage);
if (tsharp) {
tsharps.emplace_back(tsharp);
VideoCore::ImageInfo image_info{tsharp};
@ -423,7 +464,7 @@ void GraphicsPipeline::BindResources(const Liverpool::Regs& regs,
}
}
for (const auto& sampler : stage->samplers) {
auto ssharp = sampler.GetSsharp(*stage);
auto ssharp = sampler.GetSharp(*stage);
if (sampler.disable_aniso) {
const auto& tsharp = tsharps[sampler.associated_image];
if (tsharp.base_level == 0 && tsharp.last_level == 0) {

View File

@ -59,7 +59,8 @@ class GraphicsPipeline {
public:
explicit GraphicsPipeline(const Instance& instance, Scheduler& scheduler,
const GraphicsPipelineKey& key, vk::PipelineCache pipeline_cache,
std::span<const Program*, MaxShaderStages> programs);
std::span<const Shader::Info*, MaxShaderStages> stages,
std::span<const vk::ShaderModule> modules);
~GraphicsPipeline();
void BindResources(const Liverpool::Regs& regs, VideoCore::BufferCache& buffer_cache,

View File

@ -178,7 +178,7 @@ bool Instance::CreateDevice() {
return false;
}
boost::container::static_vector<const char*, 20> enabled_extensions;
boost::container::static_vector<const char*, 25> enabled_extensions;
const auto add_extension = [&](std::string_view extension) -> bool {
const auto result =
std::find_if(available_extensions.begin(), available_extensions.end(),
@ -217,6 +217,7 @@ bool Instance::CreateDevice() {
// with extensions.
tooling_info = add_extension(VK_EXT_TOOLING_INFO_EXTENSION_NAME);
const bool maintenance4 = add_extension(VK_KHR_MAINTENANCE_4_EXTENSION_NAME);
const bool maintenance5 = add_extension(VK_KHR_MAINTENANCE_5_EXTENSION_NAME);
add_extension(VK_KHR_DYNAMIC_RENDERING_EXTENSION_NAME);
add_extension(VK_EXT_SHADER_DEMOTE_TO_HELPER_INVOCATION_EXTENSION_NAME);
const bool has_sync2 = add_extension(VK_KHR_SYNCHRONIZATION_2_EXTENSION_NAME);
@ -277,6 +278,7 @@ bool Instance::CreateDevice() {
.depthBiasClamp = features.depthBiasClamp,
.multiViewport = features.multiViewport,
.samplerAnisotropy = features.samplerAnisotropy,
.vertexPipelineStoresAndAtomics = features.vertexPipelineStoresAndAtomics,
.fragmentStoresAndAtomics = features.fragmentStoresAndAtomics,
.shaderImageGatherExtended = features.shaderImageGatherExtended,
.shaderStorageImageExtendedFormats = features.shaderStorageImageExtendedFormats,
@ -299,6 +301,9 @@ bool Instance::CreateDevice() {
vk::PhysicalDeviceMaintenance4FeaturesKHR{
.maintenance4 = true,
},
vk::PhysicalDeviceMaintenance5FeaturesKHR{
.maintenance5 = true,
},
vk::PhysicalDeviceDynamicRenderingFeaturesKHR{
.dynamicRendering = true,
},
@ -344,6 +349,9 @@ bool Instance::CreateDevice() {
if (!maintenance4) {
device_chain.unlink<vk::PhysicalDeviceMaintenance4FeaturesKHR>();
}
if (!maintenance5) {
device_chain.unlink<vk::PhysicalDeviceMaintenance5FeaturesKHR>();
}
if (!custom_border_color) {
device_chain.unlink<vk::PhysicalDeviceCustomBorderColorFeaturesEXT>();
}

View File

@ -192,6 +192,11 @@ public:
return properties.limits.minStorageBufferOffsetAlignment;
}
/// Returns the minimum required alignment for texel buffers
vk::DeviceSize TexelBufferMinAlignment() const {
return properties.limits.minTexelBufferOffsetAlignment;
}
/// Returns the minimum alignemt required for accessing host-mapped device memory
vk::DeviceSize NonCoherentAtomSize() const {
return properties.limits.nonCoherentAtomSize;

View File

@ -1,147 +1,59 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "common/config.h"
#include "common/io_file.h"
#include "common/path_util.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/exception.h"
#include "shader_recompiler/recompiler.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/renderer_vulkan/renderer_vulkan.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/renderer_vulkan/vk_shader_cache.h"
extern std::unique_ptr<Vulkan::RendererVulkan> renderer;
namespace Vulkan {
using Shader::VsOutput;
[[nodiscard]] inline u64 HashCombine(const u64 seed, const u64 hash) {
return seed ^ (hash + 0x9e3779b9 + (seed << 6) + (seed >> 2));
}
void BuildVsOutputs(Shader::Info& info, const AmdGpu::Liverpool::VsOutputControl& ctl) {
const auto add_output = [&](VsOutput x, VsOutput y, VsOutput z, VsOutput w) {
if (x != VsOutput::None || y != VsOutput::None || z != VsOutput::None ||
w != VsOutput::None) {
info.vs_outputs.emplace_back(Shader::VsOutputMap{x, y, z, w});
}
};
// VS_OUT_MISC_VEC
add_output(ctl.use_vtx_point_size ? VsOutput::PointSprite : VsOutput::None,
ctl.use_vtx_edge_flag
? VsOutput::EdgeFlag
: (ctl.use_vtx_gs_cut_flag ? VsOutput::GsCutFlag : VsOutput::None),
ctl.use_vtx_kill_flag
? VsOutput::KillFlag
: (ctl.use_vtx_render_target_idx ? VsOutput::GsMrtIndex : VsOutput::None),
ctl.use_vtx_viewport_idx ? VsOutput::GsVpIndex : VsOutput::None);
// VS_OUT_CCDIST0
add_output(ctl.IsClipDistEnabled(0)
? VsOutput::ClipDist0
: (ctl.IsCullDistEnabled(0) ? VsOutput::CullDist0 : VsOutput::None),
ctl.IsClipDistEnabled(1)
? VsOutput::ClipDist1
: (ctl.IsCullDistEnabled(1) ? VsOutput::CullDist1 : VsOutput::None),
ctl.IsClipDistEnabled(2)
? VsOutput::ClipDist2
: (ctl.IsCullDistEnabled(2) ? VsOutput::CullDist2 : VsOutput::None),
ctl.IsClipDistEnabled(3)
? VsOutput::ClipDist3
: (ctl.IsCullDistEnabled(3) ? VsOutput::CullDist3 : VsOutput::None));
// VS_OUT_CCDIST1
add_output(ctl.IsClipDistEnabled(4)
? VsOutput::ClipDist4
: (ctl.IsCullDistEnabled(4) ? VsOutput::CullDist4 : VsOutput::None),
ctl.IsClipDistEnabled(5)
? VsOutput::ClipDist5
: (ctl.IsCullDistEnabled(5) ? VsOutput::CullDist5 : VsOutput::None),
ctl.IsClipDistEnabled(6)
? VsOutput::ClipDist6
: (ctl.IsCullDistEnabled(6) ? VsOutput::CullDist6 : VsOutput::None),
ctl.IsClipDistEnabled(7)
? VsOutput::ClipDist7
: (ctl.IsCullDistEnabled(7) ? VsOutput::CullDist7 : VsOutput::None));
}
Shader::Info MakeShaderInfo(Shader::Stage stage, std::span<const u32, 16> user_data,
const AmdGpu::Liverpool::Regs& regs) {
Shader::Info info{};
info.user_data = user_data;
info.stage = stage;
switch (stage) {
case Shader::Stage::Vertex: {
info.num_user_data = regs.vs_program.settings.num_user_regs;
info.num_input_vgprs = regs.vs_program.settings.vgpr_comp_cnt;
BuildVsOutputs(info, regs.vs_output_control);
break;
}
case Shader::Stage::Fragment: {
info.num_user_data = regs.ps_program.settings.num_user_regs;
for (u32 i = 0; i < regs.num_interp; i++) {
info.ps_inputs.push_back({
.param_index = regs.ps_inputs[i].input_offset.Value(),
.is_default = bool(regs.ps_inputs[i].use_default),
.is_flat = bool(regs.ps_inputs[i].flat_shade),
.default_value = regs.ps_inputs[i].default_value,
});
}
break;
}
case Shader::Stage::Compute: {
const auto& cs_pgm = regs.cs_program;
info.num_user_data = cs_pgm.settings.num_user_regs;
info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full};
info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
cs_pgm.IsTgidEnabled(2)};
info.shared_memory_size = cs_pgm.SharedMemSize();
break;
}
default:
break;
}
return info;
}
PipelineCache::PipelineCache(const Instance& instance_, Scheduler& scheduler_,
AmdGpu::Liverpool* liverpool_)
: instance{instance_}, scheduler{scheduler_}, liverpool{liverpool_}, inst_pool{8192},
block_pool{512} {
: instance{instance_}, scheduler{scheduler_}, liverpool{liverpool_},
shader_cache{std::make_unique<ShaderCache>(instance, liverpool)} {
pipeline_cache = instance.GetDevice().createPipelineCacheUnique({});
profile = Shader::Profile{
.supported_spirv = 0x00010600U,
.subgroup_size = instance.SubgroupSize(),
.support_explicit_workgroup_layout = true,
};
}
PipelineCache::~PipelineCache() = default;
const GraphicsPipeline* PipelineCache::GetGraphicsPipeline() {
const auto& regs = liverpool->regs;
// Tessellation is unsupported so skip the draw to avoid locking up the driver.
if (liverpool->regs.primitive_type == Liverpool::PrimitiveType::PatchPrimitive) {
if (regs.primitive_type == Liverpool::PrimitiveType::PatchPrimitive) {
return nullptr;
}
// There are several cases (e.g. FCE, FMask/HTile decompression) where we don't need to do an
// actual draw hence can skip pipeline creation.
if (regs.color_control.mode == Liverpool::ColorControl::OperationMode::EliminateFastClear) {
LOG_TRACE(Render_Vulkan, "FCE pass skipped");
return nullptr;
}
if (regs.color_control.mode == Liverpool::ColorControl::OperationMode::FmaskDecompress) {
// TODO: check for a valid MRT1 to promote the draw to the resolve pass.
LOG_TRACE(Render_Vulkan, "FMask decompression pass skipped");
return nullptr;
}
RefreshGraphicsKey();
const auto [it, is_new] = graphics_pipelines.try_emplace(graphics_key);
if (is_new) {
it.value() = CreateGraphicsPipeline();
it.value() = std::make_unique<GraphicsPipeline>(instance, scheduler, graphics_key,
*pipeline_cache, infos, modules);
}
const GraphicsPipeline* pipeline = it->second.get();
return pipeline;
}
const ComputePipeline* PipelineCache::GetComputePipeline() {
const auto& cs_pgm = liverpool->regs.cs_program;
ASSERT(cs_pgm.Address() != nullptr);
const auto* bininfo = Liverpool::GetBinaryInfo(cs_pgm);
compute_key = bininfo->shader_hash;
RefreshComputeKey();
const auto [it, is_new] = compute_pipelines.try_emplace(compute_key);
if (is_new) {
it.value() = CreateComputePipeline();
it.value() = std::make_unique<ComputePipeline>(instance, scheduler, *pipeline_cache,
compute_key, *infos[0], modules[0]);
}
const ComputePipeline* pipeline = it->second.get();
return pipeline;
@ -229,164 +141,37 @@ void PipelineCache::RefreshGraphicsKey() {
++remapped_cb;
}
u32 binding{};
for (u32 i = 0; i < MaxShaderStages; i++) {
if (!regs.stage_enable.IsStageEnabled(i)) {
key.stage_hashes[i] = 0;
infos[i] = nullptr;
continue;
}
auto* pgm = regs.ProgramForStage(i);
if (!pgm || !pgm->Address<u32*>()) {
key.stage_hashes[i] = 0;
infos[i] = nullptr;
continue;
}
const auto* bininfo = Liverpool::GetBinaryInfo(*pgm);
if (!bininfo->Valid()) {
key.stage_hashes[i] = 0;
infos[i] = nullptr;
continue;
}
key.stage_hashes[i] = bininfo->shader_hash;
}
}
std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline() {
const auto& regs = liverpool->regs;
// There are several cases (e.g. FCE, FMask/HTile decompression) where we don't need to do an
// actual draw hence can skip pipeline creation.
if (regs.color_control.mode == Liverpool::ColorControl::OperationMode::EliminateFastClear) {
LOG_TRACE(Render_Vulkan, "FCE pass skipped");
return {};
}
if (regs.color_control.mode == Liverpool::ColorControl::OperationMode::FmaskDecompress) {
// TODO: check for a valid MRT1 to promote the draw to the resolve pass.
LOG_TRACE(Render_Vulkan, "FMask decompression pass skipped");
return {};
}
u32 binding{};
for (u32 i = 0; i < MaxShaderStages; i++) {
if (!graphics_key.stage_hashes[i]) {
programs[i] = nullptr;
continue;
}
auto* pgm = regs.ProgramForStage(i);
const auto code = pgm->Code();
// Dump shader code if requested.
const auto stage = Shader::Stage{i};
const u64 hash = graphics_key.stage_hashes[i];
if (Config::dumpShaders()) {
DumpShader(code, hash, stage, "bin");
const GuestProgram guest_pgm{pgm, stage};
std::tie(infos[i], modules[i], key.stage_hashes[i]) =
shader_cache->GetProgram(guest_pgm, binding);
}
if (stage != Shader::Stage::Fragment && stage != Shader::Stage::Vertex) {
LOG_ERROR(Render_Vulkan, "Unsupported shader stage {}. PL creation skipped.", stage);
return {};
}
const u64 lookup_hash = HashCombine(hash, binding);
auto it = program_cache.find(lookup_hash);
if (it != program_cache.end()) {
const Program* program = it.value().get();
ASSERT(program->pgm.info.stage == stage);
programs[i] = program;
binding = program->end_binding;
continue;
}
// Recompile shader to IR.
try {
auto program = std::make_unique<Program>();
block_pool.ReleaseContents();
inst_pool.ReleaseContents();
LOG_INFO(Render_Vulkan, "Compiling {} shader {:#x}", stage, hash);
Shader::Info info = MakeShaderInfo(stage, pgm->user_data, regs);
info.pgm_base = pgm->Address<uintptr_t>();
info.pgm_hash = hash;
program->pgm =
Shader::TranslateProgram(inst_pool, block_pool, code, std::move(info), profile);
// Compile IR to SPIR-V
program->spv = Shader::Backend::SPIRV::EmitSPIRV(profile, program->pgm, binding);
if (Config::dumpShaders()) {
DumpShader(program->spv, hash, stage, "spv");
}
// Compile module and set name to hash in renderdoc
program->end_binding = binding;
program->module = CompileSPV(program->spv, instance.GetDevice());
const auto name = fmt::format("{}_{:#x}", stage, hash);
Vulkan::SetObjectName(instance.GetDevice(), program->module, name);
// Cache program
const auto [it, _] = program_cache.emplace(lookup_hash, std::move(program));
programs[i] = it.value().get();
} catch (const Shader::Exception& e) {
UNREACHABLE_MSG("{}", e.what());
}
}
return std::make_unique<GraphicsPipeline>(instance, scheduler, graphics_key, *pipeline_cache,
programs);
}
std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline() {
const auto& cs_pgm = liverpool->regs.cs_program;
const auto code = cs_pgm.Code();
// Dump shader code if requested.
if (Config::dumpShaders()) {
DumpShader(code, compute_key, Shader::Stage::Compute, "bin");
}
block_pool.ReleaseContents();
inst_pool.ReleaseContents();
// Recompile shader to IR.
try {
auto program = std::make_unique<Program>();
LOG_INFO(Render_Vulkan, "Compiling cs shader {:#x}", compute_key);
Shader::Info info =
MakeShaderInfo(Shader::Stage::Compute, cs_pgm.user_data, liverpool->regs);
info.pgm_base = cs_pgm.Address<uintptr_t>();
info.pgm_hash = compute_key;
program->pgm =
Shader::TranslateProgram(inst_pool, block_pool, code, std::move(info), profile);
// Compile IR to SPIR-V
void PipelineCache::RefreshComputeKey() {
u32 binding{};
program->spv = Shader::Backend::SPIRV::EmitSPIRV(profile, program->pgm, binding);
if (Config::dumpShaders()) {
DumpShader(program->spv, compute_key, Shader::Stage::Compute, "spv");
}
// Compile module and set name to hash in renderdoc
program->module = CompileSPV(program->spv, instance.GetDevice());
const auto name = fmt::format("cs_{:#x}", compute_key);
Vulkan::SetObjectName(instance.GetDevice(), program->module, name);
// Cache program
const auto [it, _] = program_cache.emplace(compute_key, std::move(program));
return std::make_unique<ComputePipeline>(instance, scheduler, *pipeline_cache, compute_key,
it.value().get());
} catch (const Shader::Exception& e) {
UNREACHABLE_MSG("{}", e.what());
return nullptr;
}
}
void PipelineCache::DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage,
std::string_view ext) {
using namespace Common::FS;
const auto dump_dir = GetUserPath(PathType::ShaderDir) / "dumps";
if (!std::filesystem::exists(dump_dir)) {
std::filesystem::create_directories(dump_dir);
}
const auto filename = fmt::format("{}_{:#018x}.{}", stage, hash, ext);
const auto file = IOFile{dump_dir / filename, FileAccessMode::Write};
file.WriteSpan(code);
const auto* cs_pgm = &liverpool->regs.cs_program;
const GuestProgram guest_pgm{cs_pgm, Shader::Stage::Compute};
std::tie(infos[0], modules[0], compute_key) = shader_cache->GetProgram(guest_pgm, binding);
}
} // namespace Vulkan

View File

@ -4,9 +4,6 @@
#pragma once
#include <tsl/robin_map.h>
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/ir/program.h"
#include "shader_recompiler/profile.h"
#include "video_core/renderer_vulkan/vk_compute_pipeline.h"
#include "video_core/renderer_vulkan/vk_graphics_pipeline.h"
@ -18,6 +15,7 @@ namespace Vulkan {
class Instance;
class Scheduler;
class ShaderCache;
class PipelineCache {
static constexpr size_t MaxShaderStages = 5;
@ -25,7 +23,7 @@ class PipelineCache {
public:
explicit PipelineCache(const Instance& instance, Scheduler& scheduler,
AmdGpu::Liverpool* liverpool);
~PipelineCache() = default;
~PipelineCache();
const GraphicsPipeline* GetGraphicsPipeline();
@ -33,10 +31,7 @@ public:
private:
void RefreshGraphicsKey();
void DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage, std::string_view ext);
std::unique_ptr<GraphicsPipeline> CreateGraphicsPipeline();
std::unique_ptr<ComputePipeline> CreateComputePipeline();
void RefreshComputeKey();
private:
const Instance& instance;
@ -44,15 +39,13 @@ private:
AmdGpu::Liverpool* liverpool;
vk::UniquePipelineCache pipeline_cache;
vk::UniquePipelineLayout pipeline_layout;
tsl::robin_map<size_t, std::unique_ptr<Program>> program_cache;
std::unique_ptr<ShaderCache> shader_cache;
tsl::robin_map<size_t, std::unique_ptr<ComputePipeline>> compute_pipelines;
tsl::robin_map<GraphicsPipelineKey, std::unique_ptr<GraphicsPipeline>> graphics_pipelines;
std::array<const Program*, MaxShaderStages> programs{};
Shader::Profile profile{};
std::array<const Shader::Info*, MaxShaderStages> infos{};
std::array<vk::ShaderModule, MaxShaderStages> modules{};
GraphicsPipelineKey graphics_key{};
u64 compute_key{};
Common::ObjectPool<Shader::IR::Inst> inst_pool;
Common::ObjectPool<Shader::IR::Block> block_pool;
};
} // namespace Vulkan

View File

@ -90,6 +90,45 @@ void Rasterizer::DispatchDirect() {
cmdbuf.dispatch(cs_program.dim_x, cs_program.dim_y, cs_program.dim_z);
}
void Rasterizer::DispatchIndirect(VAddr address, u32 offset, u32 size) {
RENDERER_TRACE;
const auto cmdbuf = scheduler.CommandBuffer();
const auto& cs_program = liverpool->regs.cs_program;
const ComputePipeline* pipeline = pipeline_cache.GetComputePipeline();
if (!pipeline) {
return;
}
try {
const auto has_resources = pipeline->BindResources(buffer_cache, texture_cache);
if (!has_resources) {
return;
}
} catch (...) {
UNREACHABLE();
}
scheduler.EndRendering();
cmdbuf.bindPipeline(vk::PipelineBindPoint::eCompute, pipeline->Handle());
const auto [buffer, base] = buffer_cache.ObtainBuffer(address, size, true);
const auto total_offset = base + offset;
// Emulate PFP-to-ME sync packet
const vk::BufferMemoryBarrier ib_barrier{
.srcAccessMask = vk::AccessFlagBits::eShaderWrite,
.dstAccessMask = vk::AccessFlagBits::eIndirectCommandRead,
.buffer = buffer->Handle(),
.offset = total_offset,
.size = size,
};
cmdbuf.pipelineBarrier(vk::PipelineStageFlagBits::eComputeShader,
vk::PipelineStageFlagBits::eDrawIndirect,
vk::DependencyFlagBits::eByRegion, {}, ib_barrier, {});
cmdbuf.dispatchIndirect(buffer->Handle(), total_offset);
}
u64 Rasterizer::Flush() {
const u64 current_tick = scheduler.CurrentTick();
SubmitInfo info{};

View File

@ -34,6 +34,7 @@ public:
void Draw(bool is_indexed, u32 index_offset = 0);
void DispatchDirect();
void DispatchIndirect(VAddr address, u32 offset, u32 size);
void ScopeMarkerBegin(const std::string_view& str);
void ScopeMarkerEnd();

View File

@ -0,0 +1,192 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include "common/config.h"
#include "common/io_file.h"
#include "common/path_util.h"
#include "shader_recompiler/backend/spirv/emit_spirv.h"
#include "shader_recompiler/recompiler.h"
#include "video_core/renderer_vulkan/vk_instance.h"
#include "video_core/renderer_vulkan/vk_platform.h"
#include "video_core/renderer_vulkan/vk_shader_cache.h"
#include "video_core/renderer_vulkan/vk_shader_util.h"
namespace Vulkan {
using Shader::VsOutput;
void BuildVsOutputs(Shader::Info& info, const AmdGpu::Liverpool::VsOutputControl& ctl) {
const auto add_output = [&](VsOutput x, VsOutput y, VsOutput z, VsOutput w) {
if (x != VsOutput::None || y != VsOutput::None || z != VsOutput::None ||
w != VsOutput::None) {
info.vs_outputs.emplace_back(Shader::VsOutputMap{x, y, z, w});
}
};
// VS_OUT_MISC_VEC
add_output(ctl.use_vtx_point_size ? VsOutput::PointSprite : VsOutput::None,
ctl.use_vtx_edge_flag
? VsOutput::EdgeFlag
: (ctl.use_vtx_gs_cut_flag ? VsOutput::GsCutFlag : VsOutput::None),
ctl.use_vtx_kill_flag
? VsOutput::KillFlag
: (ctl.use_vtx_render_target_idx ? VsOutput::GsMrtIndex : VsOutput::None),
ctl.use_vtx_viewport_idx ? VsOutput::GsVpIndex : VsOutput::None);
// VS_OUT_CCDIST0
add_output(ctl.IsClipDistEnabled(0)
? VsOutput::ClipDist0
: (ctl.IsCullDistEnabled(0) ? VsOutput::CullDist0 : VsOutput::None),
ctl.IsClipDistEnabled(1)
? VsOutput::ClipDist1
: (ctl.IsCullDistEnabled(1) ? VsOutput::CullDist1 : VsOutput::None),
ctl.IsClipDistEnabled(2)
? VsOutput::ClipDist2
: (ctl.IsCullDistEnabled(2) ? VsOutput::CullDist2 : VsOutput::None),
ctl.IsClipDistEnabled(3)
? VsOutput::ClipDist3
: (ctl.IsCullDistEnabled(3) ? VsOutput::CullDist3 : VsOutput::None));
// VS_OUT_CCDIST1
add_output(ctl.IsClipDistEnabled(4)
? VsOutput::ClipDist4
: (ctl.IsCullDistEnabled(4) ? VsOutput::CullDist4 : VsOutput::None),
ctl.IsClipDistEnabled(5)
? VsOutput::ClipDist5
: (ctl.IsCullDistEnabled(5) ? VsOutput::CullDist5 : VsOutput::None),
ctl.IsClipDistEnabled(6)
? VsOutput::ClipDist6
: (ctl.IsCullDistEnabled(6) ? VsOutput::CullDist6 : VsOutput::None),
ctl.IsClipDistEnabled(7)
? VsOutput::ClipDist7
: (ctl.IsCullDistEnabled(7) ? VsOutput::CullDist7 : VsOutput::None));
}
Shader::Info MakeShaderInfo(const GuestProgram& pgm, const AmdGpu::Liverpool::Regs& regs) {
Shader::Info info{};
info.user_data = pgm.user_data;
info.pgm_base = VAddr(pgm.code.data());
info.pgm_hash = pgm.hash;
info.stage = pgm.stage;
switch (pgm.stage) {
case Shader::Stage::Vertex: {
info.num_user_data = regs.vs_program.settings.num_user_regs;
info.num_input_vgprs = regs.vs_program.settings.vgpr_comp_cnt;
BuildVsOutputs(info, regs.vs_output_control);
break;
}
case Shader::Stage::Fragment: {
info.num_user_data = regs.ps_program.settings.num_user_regs;
for (u32 i = 0; i < regs.num_interp; i++) {
info.ps_inputs.push_back({
.param_index = regs.ps_inputs[i].input_offset.Value(),
.is_default = bool(regs.ps_inputs[i].use_default),
.is_flat = bool(regs.ps_inputs[i].flat_shade),
.default_value = regs.ps_inputs[i].default_value,
});
}
break;
}
case Shader::Stage::Compute: {
const auto& cs_pgm = regs.cs_program;
info.num_user_data = cs_pgm.settings.num_user_regs;
info.workgroup_size = {cs_pgm.num_thread_x.full, cs_pgm.num_thread_y.full,
cs_pgm.num_thread_z.full};
info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
cs_pgm.IsTgidEnabled(2)};
info.shared_memory_size = cs_pgm.SharedMemSize();
break;
}
default:
break;
}
return info;
}
[[nodiscard]] inline u64 HashCombine(const u64 seed, const u64 hash) {
return seed ^ (hash + 0x9e3779b9 + (seed << 6) + (seed >> 2));
}
ShaderCache::ShaderCache(const Instance& instance_, AmdGpu::Liverpool* liverpool_)
: instance{instance_}, liverpool{liverpool_}, inst_pool{8192}, block_pool{512} {
profile = Shader::Profile{
.supported_spirv = 0x00010600U,
.subgroup_size = instance.SubgroupSize(),
.support_explicit_workgroup_layout = true,
};
}
vk::ShaderModule ShaderCache::CompileModule(Shader::Info& info, std::span<const u32> code,
size_t perm_idx, u32& binding) {
LOG_INFO(Render_Vulkan, "Compiling {} shader {:#x} {}", info.stage, info.pgm_hash,
perm_idx != 0 ? "(permutation)" : "");
if (Config::dumpShaders()) {
DumpShader(code, info.pgm_hash, info.stage, perm_idx, "bin");
}
block_pool.ReleaseContents();
inst_pool.ReleaseContents();
const auto ir_program = Shader::TranslateProgram(inst_pool, block_pool, code, info, profile);
// Compile IR to SPIR-V
const auto spv = Shader::Backend::SPIRV::EmitSPIRV(profile, ir_program, binding);
if (Config::dumpShaders()) {
DumpShader(spv, info.pgm_hash, info.stage, perm_idx, "spv");
}
// Create module and set name to hash in renderdoc
const auto module = CompileSPV(spv, instance.GetDevice());
ASSERT(module != VK_NULL_HANDLE);
const auto name = fmt::format("{}_{:#x}_{}", info.stage, info.pgm_hash, perm_idx);
Vulkan::SetObjectName(instance.GetDevice(), module, name);
return module;
}
Program* ShaderCache::CreateProgram(const GuestProgram& pgm, u32& binding) {
Program* program = program_pool.Create(MakeShaderInfo(pgm, liverpool->regs));
u32 start_binding = binding;
const auto module = CompileModule(program->info, pgm.code, 0, binding);
program->modules.emplace_back(module, StageSpecialization{program->info, start_binding});
return program;
}
std::tuple<const Shader::Info*, vk::ShaderModule, u64> ShaderCache::GetProgram(
const GuestProgram& pgm, u32& binding) {
auto [it_pgm, new_program] = program_cache.try_emplace(pgm.hash);
if (new_program) {
auto program = CreateProgram(pgm, binding);
const auto module = program->modules.back().module;
it_pgm.value() = program;
return std::make_tuple(&program->info, module, HashCombine(pgm.hash, 0));
}
Program* program = it_pgm->second;
const auto& info = program->info;
size_t perm_idx = program->modules.size();
StageSpecialization spec{info, binding};
vk::ShaderModule module{};
const auto it = std::ranges::find(program->modules, spec, &Program::Module::spec);
if (it == program->modules.end()) {
auto new_info = MakeShaderInfo(pgm, liverpool->regs);
module = CompileModule(new_info, pgm.code, perm_idx, binding);
program->modules.emplace_back(module, std::move(spec));
} else {
binding += info.NumBindings();
module = it->module;
perm_idx = std::distance(program->modules.begin(), it);
}
return std::make_tuple(&info, module, HashCombine(pgm.hash, perm_idx));
}
void ShaderCache::DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage,
size_t perm_idx, std::string_view ext) {
using namespace Common::FS;
const auto dump_dir = GetUserPath(PathType::ShaderDir) / "dumps";
if (!std::filesystem::exists(dump_dir)) {
std::filesystem::create_directories(dump_dir);
}
const auto filename = fmt::format("{}_{:#018x}_{}.{}", stage, hash, perm_idx, ext);
const auto file = IOFile{dump_dir / filename, FileAccessMode::Write};
file.WriteSpan(code);
}
} // namespace Vulkan

View File

@ -0,0 +1,156 @@
// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <bitset>
#include <boost/container/small_vector.hpp>
#include <tsl/robin_map.h>
#include "common/object_pool.h"
#include "shader_recompiler/ir/basic_block.h"
#include "shader_recompiler/profile.h"
#include "shader_recompiler/runtime_info.h"
#include "video_core/amdgpu/liverpool.h"
#include "video_core/renderer_vulkan/vk_common.h"
namespace Vulkan {
class Instance;
struct BufferSpecialization {
u16 stride : 14;
u16 is_storage : 1;
auto operator<=>(const BufferSpecialization&) const = default;
};
struct TextureBufferSpecialization {
bool is_integer;
auto operator<=>(const TextureBufferSpecialization&) const = default;
};
struct ImageSpecialization {
AmdGpu::ImageType type;
bool is_integer;
auto operator<=>(const ImageSpecialization&) const = default;
};
struct StageSpecialization {
static constexpr size_t MaxStageResources = 32;
const Shader::Info* info;
std::bitset<MaxStageResources> bitset{};
boost::container::small_vector<BufferSpecialization, 16> buffers;
boost::container::small_vector<TextureBufferSpecialization, 8> tex_buffers;
boost::container::small_vector<ImageSpecialization, 8> images;
u32 start_binding{};
void ForEachSharp(u32& binding, auto& spec_list, auto& desc_list, auto&& func) {
for (const auto& desc : desc_list) {
auto& spec = spec_list.emplace_back();
const auto sharp = desc.GetSharp(*info);
if (!sharp) {
binding++;
continue;
}
bitset.set(binding++);
func(spec, desc, sharp);
}
}
StageSpecialization(const Shader::Info& info_, u32 start_binding_)
: info{&info_}, start_binding{start_binding_} {
u32 binding{};
ForEachSharp(binding, buffers, info->buffers,
[](auto& spec, const auto& desc, AmdGpu::Buffer sharp) {
spec.stride = sharp.GetStride();
spec.is_storage = desc.IsStorage(sharp);
});
ForEachSharp(binding, tex_buffers, info->texture_buffers,
[](auto& spec, const auto& desc, AmdGpu::Buffer sharp) {
spec.is_integer = AmdGpu::IsInteger(sharp.GetNumberFmt());
});
ForEachSharp(binding, images, info->images,
[](auto& spec, const auto& desc, AmdGpu::Image sharp) {
spec.type = sharp.GetType();
spec.is_integer = AmdGpu::IsInteger(sharp.GetNumberFmt());
});
}
bool operator==(const StageSpecialization& other) const {
if (start_binding != other.start_binding) {
return false;
}
u32 binding{};
for (u32 i = 0; i < buffers.size(); i++) {
if (other.bitset[binding++] && buffers[i] != other.buffers[i]) {
return false;
}
}
for (u32 i = 0; i < tex_buffers.size(); i++) {
if (other.bitset[binding++] && tex_buffers[i] != other.tex_buffers[i]) {
return false;
}
}
for (u32 i = 0; i < images.size(); i++) {
if (other.bitset[binding++] && images[i] != other.images[i]) {
return false;
}
}
return true;
}
};
struct Program {
struct Module {
vk::ShaderModule module;
StageSpecialization spec;
};
Shader::Info info;
boost::container::small_vector<Module, 8> modules;
explicit Program(const Shader::Info& info_) : info{info_} {}
};
struct GuestProgram {
Shader::Stage stage;
std::span<const u32, AmdGpu::Liverpool::NumShaderUserData> user_data;
std::span<const u32> code;
u64 hash;
explicit GuestProgram(const auto* pgm, Shader::Stage stage_)
: stage{stage_}, user_data{pgm->user_data}, code{pgm->Code()} {
const auto* bininfo = AmdGpu::Liverpool::GetBinaryInfo(*pgm);
hash = bininfo->shader_hash;
}
};
class ShaderCache {
public:
explicit ShaderCache(const Instance& instance, AmdGpu::Liverpool* liverpool);
~ShaderCache() = default;
std::tuple<const Shader::Info*, vk::ShaderModule, u64> GetProgram(const GuestProgram& pgm,
u32& binding);
private:
void DumpShader(std::span<const u32> code, u64 hash, Shader::Stage stage, size_t perm_idx,
std::string_view ext);
vk::ShaderModule CompileModule(Shader::Info& info, std::span<const u32> code, size_t perm_idx,
u32& binding);
Program* CreateProgram(const GuestProgram& pgm, u32& binding);
private:
const Instance& instance;
AmdGpu::Liverpool* liverpool;
Shader::Profile profile{};
tsl::robin_map<size_t, Program*> program_cache;
Common::ObjectPool<Shader::IR::Inst> inst_pool;
Common::ObjectPool<Shader::IR::Block> block_pool;
Common::ObjectPool<Program> program_pool;
};
} // namespace Vulkan

View File

@ -50,9 +50,9 @@ vk::ComponentSwizzle ConvertComponentSwizzle(u32 dst_sel) {
}
bool IsIdentityMapping(u32 dst_sel, u32 num_components) {
return (num_components == 1 && dst_sel == 0b100) ||
(num_components == 2 && dst_sel == 0b101'100) ||
(num_components == 3 && dst_sel == 0b110'101'100) ||
return (num_components == 1 && dst_sel == 0b001'000'000'100) ||
(num_components == 2 && dst_sel == 0b001'000'101'100) ||
(num_components == 3 && dst_sel == 0b001'110'101'100) ||
(num_components == 4 && dst_sel == 0b111'110'101'100);
}

View File

@ -34,10 +34,10 @@ TextureCache::TextureCache(const Vulkan::Instance& instance_, Vulkan::Scheduler&
TextureCache::~TextureCache() = default;
void TextureCache::InvalidateMemory(VAddr address, size_t size, bool from_compute) {
void TextureCache::InvalidateMemory(VAddr address, size_t size) {
std::unique_lock lock{mutex};
ForEachImageInRegion(address, size, [&](ImageId image_id, Image& image) {
if (from_compute && !image.Overlaps(address, size)) {
if (!image.Overlaps(address, size)) {
return;
}
// Ensure image is reuploaded when accessed again.

View File

@ -38,7 +38,7 @@ public:
~TextureCache();
/// Invalidates any image in the logical page range.
void InvalidateMemory(VAddr address, size_t size, bool from_compute = false);
void InvalidateMemory(VAddr address, size_t size);
/// Evicts any images that overlap the unmapped range.
void UnmapMemory(VAddr cpu_addr, size_t size);

View File

@ -187,6 +187,7 @@ vk::Format DemoteImageFormatForDetiling(vk::Format format) {
case vk::Format::eR32Uint:
case vk::Format::eR16G16Sfloat:
case vk::Format::eR16G16Unorm:
case vk::Format::eB10G11R11UfloatPack32:
return vk::Format::eR32Uint;
case vk::Format::eBc1RgbaSrgbBlock:
case vk::Format::eBc1RgbaUnormBlock:
@ -202,6 +203,7 @@ vk::Format DemoteImageFormatForDetiling(vk::Format format) {
case vk::Format::eBc3SrgbBlock:
case vk::Format::eBc3UnormBlock:
case vk::Format::eBc5UnormBlock:
case vk::Format::eBc5SnormBlock:
case vk::Format::eBc7SrgbBlock:
case vk::Format::eBc7UnormBlock:
case vk::Format::eBc6HUfloatBlock: