diff --git a/.gitmodules b/.gitmodules index 94996586..95b0fc0b 100644 --- a/.gitmodules +++ b/.gitmodules @@ -81,4 +81,7 @@ [submodule "externals/ffmpeg-core"] path = externals/ffmpeg-core url = https://github.com/shadps4-emu/ext-ffmpeg-core.git - shallow = true \ No newline at end of file + shallow = true +[submodule "externals/half"] + path = externals/half + url = https://github.com/ROCm/half.git diff --git a/CMakeLists.txt b/CMakeLists.txt index dfc6528d..8c7e03e3 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -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) diff --git a/externals/CMakeLists.txt b/externals/CMakeLists.txt index 6fe73a29..de0317ff 100644 --- a/externals/CMakeLists.txt +++ b/externals/CMakeLists.txt @@ -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) - option(BUILD_TZ_LIB "" ON) - option(USE_SYSTEM_TZ_DB "" ON) - add_subdirectory(date) +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 diff --git a/externals/half b/externals/half new file mode 160000 index 00000000..1ddada22 --- /dev/null +++ b/externals/half @@ -0,0 +1 @@ +Subproject commit 1ddada225144cac0de8f6b5c0dd9acffd99a2e68 diff --git a/src/common/path_util.cpp b/src/common/path_util.cpp index 5d5c9eba..8d369fc7 100644 --- a/src/common/path_util.cpp +++ b/src/common/path_util.cpp @@ -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; }(); diff --git a/src/common/path_util.h b/src/common/path_util.h index 8922de9f..bee93c1b 100644 --- a/src/common/path_util.h +++ b/src/common/path_util.h @@ -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"; diff --git a/src/core/cpu_patches.cpp b/src/core/cpu_patches.cpp index 42318822..55bbf23b 100644 --- a/src/core/cpu_patches.cpp +++ b/src/core/cpu_patches.cpp @@ -15,6 +15,7 @@ #else #include #ifdef __APPLE__ +#include #include #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(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(operand.type)); + return operand.imm.value.u; +} + static std::unique_ptr 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(®ister_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); } } +/// 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[reinterpret_cast(stack_pointer_slot * sizeof(void*))], rsp); + c.putSeg(gs); + c.mov(rsp, qword[reinterpret_cast(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(patch_stack_slot * sizeof(void*))], rsp); + c.putSeg(gs); + c.mov(rsp, qword[reinterpret_cast(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 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; + SaveStack(c); for (const auto& reg : regs) { - const auto offset = reinterpret_cast(register_save_slots[index++] * sizeof(void*)); - - c.putSeg(gs); - c.mov(qword[offset], reg.cvt64()); + 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 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(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(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(src.get())); + } else { + c.lea(rsi, src->getAddress()); + } + + // Load float count into third parameter. + c.mov(rdx, float_count); + + c.mov(rax, reinterpret_cast(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(&dst), ptr[rsp]); + } else { + c.movdqu(*reinterpret_cast(&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_cast, + half_float::half_cast, + half_float::half_cast, +}; + +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(&src)); + } else { + c.movdqu(ptr[rsp], *reinterpret_cast(&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(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(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 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 }; diff --git a/src/core/cpu_patches.h b/src/core/cpu_patches.h index 45adbeda..9126074e 100644 --- a/src/core/cpu_patches.h +++ b/src/core/cpu_patches.h @@ -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); diff --git a/src/core/file_format/pkg.cpp b/src/core/file_format/pkg.cpp index fe3c3454..d86f3b28 100644 --- a/src/core/file_format/pkg.cpp +++ b/src/core/file_format/pkg.cpp @@ -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, 7> key1; std::array 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(pkg.data() + entry.offset, entry.size); + file.Seek(entry.offset); + + std::vector data; + data.resize(entry.size); + file.ReadRaw(data.data(), entry.size); + out.WriteRaw(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(pkg.data() + entry.offset, entry.size); + file.Seek(entry.offset); + + std::vector data; + data.resize(entry.size); + file.ReadRaw(data.data(), entry.size); + out.WriteRaw(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 cipherNp(pkg.data() + entry.offset, entry.size); + file.Seek(entry.offset); + + std::vector data; + data.resize(entry.size); + file.ReadRaw(data.data(), entry.size); + + std::span cipherNp(data.data(), entry.size); std::array 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,28 +244,31 @@ 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. - // Read encrypted pfs_image - std::vector pfs_encrypted(length); - file.Seek(pkgheader.pfs_image_offset); - file.Read(pfs_encrypted); - file.Close(); - // Decrypt the pfs_image. - std::vector pfs_decrypted(length); - PKG::crypto.decryptPFS(dataKey, tweakKey, pfs_encrypted, pfs_decrypted, 0); - - // Retrieve PFSC from decrypted pfs_image. - pfsc_offset = GetPFSCOffset(pfs_decrypted); + int num_blocks = 0; std::vector pfsc(length); - std::memcpy(pfsc.data(), pfs_decrypted.data() + pfsc_offset, length - pfsc_offset); + if (length != 0) { + // Read encrypted pfs_image + std::vector pfs_encrypted(length); + file.Seek(pkgheader.pfs_image_offset); + file.Read(pfs_encrypted); + file.Close(); + // Decrypt the pfs_image. + std::vector pfs_decrypted(length); + PKG::crypto.decryptPFS(dataKey, tweakKey, pfs_encrypted, pfs_decrypted, 0); - PFSCHdr pfsChdr; - std::memcpy(&pfsChdr, pfsc.data(), sizeof(pfsChdr)); + // Retrieve PFSC from decrypted pfs_image. + pfsc_offset = GetPFSCOffset(pfs_decrypted); + std::memcpy(pfsc.data(), pfs_decrypted.data() + pfsc_offset, length - pfsc_offset); - const int num_blocks = (int)(pfsChdr.data_length / pfsChdr.block_sz2); - sectorMap.resize(num_blocks + 1); // 8 bytes, need extra 1 to get the last offset. + PFSCHdr pfsChdr; + std::memcpy(&pfsChdr, pfsc.data(), sizeof(pfsChdr)); - for (int i = 0; i < num_blocks + 1; i++) { - std::memcpy(§orMap[i], pfsc.data() + pfsChdr.block_offsets + i * 8, 8); + 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(§orMap[i], pfsc.data() + pfsChdr.block_offsets + i * 8, 8); + } } u32 ent_size = 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; } diff --git a/src/core/file_format/pkg.h b/src/core/file_format/pkg.h index b6b09a19..d30d50b4 100644 --- a/src/core/file_format/pkg.h +++ b/src/core/file_format/pkg.h @@ -149,7 +149,6 @@ public: private: Crypto crypto; TRP trp; - std::vector pkg; u64 pkgSize = 0; char pkgTitleID[9]; PKGHeader pkgheader; diff --git a/src/core/libraries/app_content/app_content.cpp b/src/core/libraries/app_content/app_content.cpp index ab3c901a..c2523124 100644 --- a/src/core/libraries/app_content/app_content.cpp +++ b/src/core/libraries/app_content/app_content.cpp @@ -7,14 +7,33 @@ #include #include #include + #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 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"); - return ORBIS_OK; +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::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"); - return ORBIS_OK; +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"); - return ORBIS_OK; +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::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; } @@ -324,4 +441,4 @@ void RegisterlibSceAppContent(Core::Loader::SymbolsResolver* sym) { sceAppContentGetDownloadedStoreCountry); }; -} // namespace Libraries::AppContent \ No newline at end of file +} // namespace Libraries::AppContent diff --git a/src/core/libraries/app_content/app_content.h b/src/core/libraries/app_content/app_content.h index 3e6f9b54..a16da5b4 100644 --- a/src/core/libraries/app_content/app_content.h +++ b/src/core/libraries/app_content/app_content.h @@ -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); diff --git a/src/core/libraries/avplayer/avplayer.cpp b/src/core/libraries/avplayer/avplayer.cpp index 406583a8..23e1e987 100644 --- a/src/core/libraries/avplayer/avplayer.cpp +++ b/src/core/libraries/avplayer/avplayer.cpp @@ -325,4 +325,4 @@ void RegisterlibSceAvPlayer(Core::Loader::SymbolsResolver* sym) { LIB_FUNCTION("yN7Jhuv8g24", "libSceAvPlayer", 1, "libSceAvPlayer", 1, 0, sceAvPlayerVprintf); }; -} // namespace Libraries::AvPlayer \ No newline at end of file +} // namespace Libraries::AvPlayer diff --git a/src/core/libraries/error_codes.h b/src/core/libraries/error_codes.h index 1453b0a8..094ea660 100644 --- a/src/core/libraries/error_codes.h +++ b/src/core/libraries/error_codes.h @@ -460,4 +460,6 @@ constexpr int ORBIS_AVPLAYER_ERROR_INFO_AES_ENCRY = 0x806A00B5; constexpr int ORBIS_AVPLAYER_ERROR_INFO_OTHER_ENCRY = 0x806A00BF; // AppContent library -constexpr int ORBIS_APP_CONTENT_ERROR_PARAMETER = 0x80D90002; \ No newline at end of file +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; \ No newline at end of file diff --git a/src/core/libraries/kernel/thread_management.cpp b/src/core/libraries/kernel/thread_management.cpp index 605d0d29..a2befd4c 100644 --- a/src/core/libraries/kernel/thread_management.cpp +++ b/src/core/libraries/kernel/thread_management.cpp @@ -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(arg); Common::SetCurrentThreadName(thread->name.c_str()); auto* linker = Common::Singleton::Instance(); + Core::InitializeThreadPatchStack(); linker->InitTlsForThread(false); void* ret = nullptr; g_pthread_self = thread; diff --git a/src/core/libraries/kernel/threads/semaphore.cpp b/src/core/libraries/kernel/threads/semaphore.cpp index 5304dc57..e2f43803 100644 --- a/src/core/libraries/kernel/threads/semaphore.cpp +++ b/src/core/libraries/kernel/threads/semaphore.cpp @@ -2,9 +2,8 @@ // SPDX-License-Identifier: GPL-2.0-or-later #include +#include #include -#include -#include #include #include "common/assert.h" #include "common/logging/log.h" @@ -13,9 +12,6 @@ namespace Libraries::Kernel { -using ListBaseHook = - boost::intrusive::list_base_hook>; - 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, - boost::intrusive::constant_time_size>; - WaitingThreads wait_list; + std::list wait_list; std::string name; std::atomic token_count; std::mutex mutex; diff --git a/src/core/linker.cpp b/src/core/linker.cpp index 9783ad96..4ef62c4a 100644 --- a/src/core/linker.cpp +++ b/src/core/linker.cpp @@ -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) { diff --git a/src/emulator.cpp b/src/emulator.cpp index b12bb859..85a4d745 100644 --- a/src/emulator.cpp +++ b/src/emulator.cpp @@ -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::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; diff --git a/src/qt_gui/cheats_patches.cpp b/src/qt_gui/cheats_patches.cpp index 662d52cc..fd0bc4e1 100644 --- a/src/qt_gui/cheats_patches.cpp +++ b/src/qt_gui/cheats_patches.cpp @@ -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); @@ -1261,4 +1262,4 @@ void CheatsPatches::onPatchCheckBoxHovered(QCheckBox* checkBox, bool hovered) { } else { instructionsTextEdit->setText(defaultTextEdit); } -} \ No newline at end of file +} diff --git a/src/qt_gui/cheats_patches.h b/src/qt_gui/cheats_patches.h index 7a68829c..a9932886 100644 --- a/src/qt_gui/cheats_patches.h +++ b/src/qt_gui/cheats_patches.h @@ -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); diff --git a/src/qt_gui/main_window.cpp b/src/qt_gui/main_window.cpp index 23668ef7..944ff4a6 100644 --- a/src/qt_gui/main_window.cpp +++ b/src/qt_gui/main_window.cpp @@ -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")); + + 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")) { - psf.open("", pkg.sfo); 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 { + 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%1\nWould you like to overwrite?")) - .arg(QString::fromStdString(extract_path.string()))); + 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,45 +725,47 @@ void MainWindow::InstallDragDropPkg(std::filesystem::path file, int pkgNum, int } else { int nfiles = pkg.GetNumberOfFiles(); - QVector indices; - for (int i = 0; i < nfiles; i++) { - indices.append(i); - } - - QProgressDialog dialog; - dialog.setWindowTitle(tr("PKG Extraction")); - dialog.setWindowModality(Qt::WindowModal); - QString extractmsg = QString(tr("Extracting PKG %1/%2")).arg(pkgNum).arg(nPkg); - dialog.setLabelText(extractmsg); - dialog.setAutoClose(true); - dialog.setRange(0, nfiles); - - QFutureWatcher futureWatcher; - connect(&futureWatcher, &QFutureWatcher::finished, this, [=, this]() { - if (pkgNum == nPkg) { - QString path = QString::fromStdString(Config::getGameInstallDir()); - QMessageBox extractMsgBox(this); - extractMsgBox.setWindowTitle(tr("Extraction Finished")); - extractMsgBox.setText( - QString(tr("Game successfully installed at %1")).arg(path)); - extractMsgBox.addButton(QMessageBox::Ok); - extractMsgBox.setDefaultButton(QMessageBox::Ok); - connect(&extractMsgBox, &QMessageBox::buttonClicked, this, - [&](QAbstractButton* button) { - if (extractMsgBox.button(QMessageBox::Ok) == button) { - extractMsgBox.close(); - emit ExtractionFinished(); - } - }); - extractMsgBox.exec(); + if (nfiles > 0) { + QVector indices; + for (int i = 0; i < nfiles; i++) { + indices.append(i); } - }); - connect(&dialog, &QProgressDialog::canceled, [&]() { futureWatcher.cancel(); }); - connect(&futureWatcher, &QFutureWatcher::progressValueChanged, &dialog, - &QProgressDialog::setValue); - futureWatcher.setFuture( - QtConcurrent::map(indices, [&](int index) { pkg.ExtractFiles(index); })); - dialog.exec(); + + QProgressDialog dialog; + dialog.setWindowTitle(tr("PKG Extraction")); + dialog.setWindowModality(Qt::WindowModal); + QString extractmsg = QString(tr("Extracting PKG %1/%2")).arg(pkgNum).arg(nPkg); + dialog.setLabelText(extractmsg); + dialog.setAutoClose(true); + dialog.setRange(0, nfiles); + + QFutureWatcher futureWatcher; + connect(&futureWatcher, &QFutureWatcher::finished, this, [=, this]() { + if (pkgNum == nPkg) { + QString path = QString::fromStdString(Config::getGameInstallDir()); + QMessageBox extractMsgBox(this); + extractMsgBox.setWindowTitle(tr("Extraction Finished")); + extractMsgBox.setText( + QString(tr("Game successfully installed at %1")).arg(path)); + extractMsgBox.addButton(QMessageBox::Ok); + extractMsgBox.setDefaultButton(QMessageBox::Ok); + connect(&extractMsgBox, &QMessageBox::buttonClicked, this, + [&](QAbstractButton* button) { + if (extractMsgBox.button(QMessageBox::Ok) == button) { + extractMsgBox.close(); + emit ExtractionFinished(); + } + }); + extractMsgBox.exec(); + } + }); + connect(&dialog, &QProgressDialog::canceled, [&]() { futureWatcher.cancel(); }); + connect(&futureWatcher, &QFutureWatcher::progressValueChanged, &dialog, + &QProgressDialog::setValue); + futureWatcher.setFuture( + QtConcurrent::map(indices, [&](int index) { pkg.ExtractFiles(index); })); + dialog.exec(); + } } } else { QMessageBox::critical(this, tr("PKG ERROR"), diff --git a/src/qt_gui/translations/da_DK.ts b/src/qt_gui/translations/da_DK.ts index de9097f2..c67d29b1 100644 --- a/src/qt_gui/translations/da_DK.ts +++ b/src/qt_gui/translations/da_DK.ts @@ -566,44 +566,54 @@ PKG-udtrækning - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch opdaget!\nPKG- og spilversioner stemmer overens!: %1\nVil du + + Patch detected! + Opdatering detekteret! - - to overwrite? - overskrive? + + PKG and Game versions match: + PKG og spilversioner matcher: - - Patch detected!\nPKG Version %1 is older - Patch opdaget!\nPKG-version %1 er ældre + + Would you like to overwrite? + Vil du overskrive? - - than installed version!: %2\nWould you like - end installeret version!: %2\nVil du + + PKG Version %1 is older than installed version: + PKG Version %1 er ældre end den installerede version: - - to overwrite? - overskrive? + + Game is installed: + Spillet er installeret: - - Patch detected!\nGame is installed: %1\nWould you like - Patch opdaget!\nSpillet er installeret: %1\nVil du + + Would you like to install Patch: + Vil du installere opdateringen: - - to install Patch: %2? - installere patch: %2? + + DLC Installation + DLC Installation - - Game already installed\n%1\nWould you like to overwrite? - Spil allerede installeret\n%1\nVil du overskrive? + + Would you like to install DLC: %1? + Vil du installere DLC: %1? + + + + DLC already installed: + DLC allerede installeret: + + + + Game already installed + Spillet er allerede installeret diff --git a/src/qt_gui/translations/de.ts b/src/qt_gui/translations/de.ts index 6183c812..c833f2e2 100644 --- a/src/qt_gui/translations/de.ts +++ b/src/qt_gui/translations/de.ts @@ -566,44 +566,54 @@ PKG-Extraktion - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch erkannt!\nPKG- und Spielversion stimmen überein!: %1\nMöchten Sie + + Patch detected! + Patch erkannt! - - to overwrite? - überschreiben? + + PKG and Game versions match: + PKG- und Spielversionen stimmen überein: - - Patch detected!\nPKG Version %1 is older - Patch erkannt!\nPKG-Version %1 ist älter + + Would you like to overwrite? + Würden Sie gerne überschreiben? - - than installed version!: %2\nWould you like - als die installierte Version!: %2\nMöchten Sie + + PKG Version %1 is older than installed version: + PKG-Version %1 ist älter als die installierte Version: - - to overwrite? - überschreiben? + + Game is installed: + Spiel ist installiert: - - Patch detected!\nGame is installed: %1\nWould you like - Patch erkannt!\nSpiel ist installiert: %1\nMöchten Sie + + Would you like to install Patch: + Möchten Sie den Patch installieren: - - to install Patch: %2? - Patch installieren: %2? + + DLC Installation + DLC-Installation - - Game already installed\n%1\nWould you like to overwrite? - Spiel bereits installiert\n%1\nMöchten Sie überschreiben? + + Would you like to install DLC: %1? + Würden Sie gerne DLC installieren: %1? + + + + DLC already installed: + DLC bereits installiert: + + + + Game already installed + Spiel bereits installiert diff --git a/src/qt_gui/translations/el.ts b/src/qt_gui/translations/el.ts index 5d4a15d8..ef831fb0 100644 --- a/src/qt_gui/translations/el.ts +++ b/src/qt_gui/translations/el.ts @@ -566,44 +566,54 @@ Εξαγωγή PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Ανίχνευση Patch!\nΟι εκδόσεις PKG και παιχνιδιού ταιριάζουν!: %1\nΘέλετε + + Patch detected! + Αναγνώριση ενημέρωσης! - - to overwrite? - να αντικαταστήσετε; + + PKG and Game versions match: + Οι εκδόσεις PKG και παιχνιδιού ταιριάζουν: - - Patch detected!\nPKG Version %1 is older - Ανίχνευση Patch!\nΗ έκδοση PKG %1 είναι παλαιότερη + + Would you like to overwrite? + Θέλετε να αντικαταστήσετε; - - than installed version!: %2\nWould you like - από την εγκατεστημένη έκδοση!: %2\nΘέλετε + + PKG Version %1 is older than installed version: + Η έκδοση PKG %1 είναι παλαιότερη από την εγκατεστημένη έκδοση: - - to overwrite? - να αντικαταστήσετε; + + Game is installed: + Το παιχνίδι είναι εγκατεστημένο: - - Patch detected!\nGame is installed: %1\nWould you like - Ανίχνευση Patch!\nΤο παιχνίδι είναι εγκατεστημένο: %1\nΘέλετε + + Would you like to install Patch: + Θέλετε να εγκαταστήσετε την ενημέρωση: - - to install Patch: %2? - να εγκαταστήσετε το Patch: %2; + + DLC Installation + Εγκατάσταση DLC - - Game already installed\n%1\nWould you like to overwrite? - Το παιχνίδι είναι ήδη εγκατεστημένο\n%1\nΘέλετε να αντικαταστήσετε; + + Would you like to install DLC: %1? + Θέλετε να εγκαταστήσετε το DLC: %1; + + + + DLC already installed: + DLC ήδη εγκατεστημένο: + + + + Game already installed + Παιχνίδι ήδη εγκατεστημένο diff --git a/src/qt_gui/translations/en.ts b/src/qt_gui/translations/en.ts index aa8a1a5d..b3c3b699 100644 --- a/src/qt_gui/translations/en.ts +++ b/src/qt_gui/translations/en.ts @@ -566,44 +566,54 @@ PKG Extraction - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch detected!\nPKG and Game versions match!: %1\nWould you like - - - - to overwrite? - to overwrite? - - - - Patch detected!\nPKG Version %1 is older - Patch detected!\nPKG Version %1 is older - - - - than installed version!: %2\nWould you like - than installed version!: %2\nWould you like - - - - to overwrite? - to overwrite? + + Patch detected! + Patch detected! - Patch detected!\nGame is installed: %1\nWould you like - Patch detected!\nGame is installed: %1\nWould you like + PKG and Game versions match: + PKG and Game versions match: - to install Patch: %2? - to install Patch: %2? + Would you like to overwrite? + Would you like to overwrite? - - Game already installed\n%1\nWould you like to overwrite? - Game already installed\n%1\nWould you like to overwrite? + + PKG Version %1 is older than installed version: + PKG Version %1 is older than installed version: + + + + Game is installed: + Game is installed: + + + + Would you like to install Patch: + Would you like to install Patch: + + + + DLC Installation + DLC Installation + + + + Would you like to install DLC: %1? + Would you like to install DLC: %1? + + + + DLC already installed: + DLC already installed: + + + + Game already installed + Game already installed diff --git a/src/qt_gui/translations/es_ES.ts b/src/qt_gui/translations/es_ES.ts index 5b7ad4b6..c34dc3d4 100644 --- a/src/qt_gui/translations/es_ES.ts +++ b/src/qt_gui/translations/es_ES.ts @@ -566,44 +566,54 @@ Extracción de PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - ¡Parche detectado!\n¡La versión de PKG y del juego coinciden!: %1\n¿Te gustaría - - - - to overwrite? - ¿sobrescribir? - - - - Patch detected!\nPKG Version %1 is older - ¡Parche detectado!\nLa versión de PKG %1 es más antigua - - - - than installed version!: %2\nWould you like - que la versión instalada!: %2\n¿Te gustaría - - - - to overwrite? - ¿sobrescribir? + + Patch detected! + ¡Actualización detectada! - Patch detected!\nGame is installed: %1\nWould you like - ¡Parche detectado!\nJuego está instalado: %1\n¿Te gustaría + PKG and Game versions match: + Las versiones de PKG y del juego coinciden: - to install Patch: %2? - ¿instalar el parche: %2? + Would you like to overwrite? + ¿Desea sobrescribir? - - Game already installed\n%1\nWould you like to overwrite? - Juego ya instalado\n%1\n¿Te gustaría sobrescribirlo? + + PKG Version %1 is older than installed version: + La versión de PKG %1 es más antigua que la versión instalada: + + + + Game is installed: + El juego está instalado: + + + + Would you like to install Patch: + ¿Desea instalar la actualización: + + + + DLC Installation + Instalación de DLC + + + + Would you like to install DLC: %1? + ¿Desea instalar el DLC: %1? + + + + DLC already installed: + DLC ya instalado: + + + + Game already installed + Juego ya instalado diff --git a/src/qt_gui/translations/fi.ts b/src/qt_gui/translations/fi.ts index 70345efb..d667dd37 100644 --- a/src/qt_gui/translations/fi.ts +++ b/src/qt_gui/translations/fi.ts @@ -566,44 +566,54 @@ PKG:n purku - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Korjaus havaittu!\nPKG:n ja pelin versiot vastaavat!: %1\nHaluatko + + Patch detected! + Päivitys havaittu! - - to overwrite? - korvata? + + PKG and Game versions match: + PKG- ja peliversiot vastaavat: - - Patch detected!\nPKG Version %1 is older - Korjaus havaittu!\nPKG Version %1 on vanhempi + + Would you like to overwrite? + Haluatko korvata? - - than installed version!: %2\nWould you like - kuin asennettu versio!: %2\nHaluatko + + PKG Version %1 is older than installed version: + PKG-versio %1 on vanhempi kuin asennettu versio: - - to overwrite? - korvata? + + Game is installed: + Peli on asennettu: - - Patch detected!\nGame is installed: %1\nWould you like - Korjaus havaittu!\nPeli on asennettu: %1\nHaluatko + + Would you like to install Patch: + Haluatko asentaa päivityksen: - - to install Patch: %2? - asentaa korjaus: %2? + + DLC Installation + DLC-asennus - - Game already installed\n%1\nWould you like to overwrite? - Peli on jo asennettu\n%1\nHaluatko korvata sen? + + Would you like to install DLC: %1? + Haluatko asentaa DLC:n: %1? + + + + DLC already installed: + DLC on jo asennettu: + + + + Game already installed + Peli on jo asennettu diff --git a/src/qt_gui/translations/fr.ts b/src/qt_gui/translations/fr.ts index 371da961..388912d2 100644 --- a/src/qt_gui/translations/fr.ts +++ b/src/qt_gui/translations/fr.ts @@ -566,44 +566,54 @@ Extraction du PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch détecté !\nLa version du PKG et du jeu correspondent : %1\nSouhaitez-vous + + Patch detected! + Patch détecté ! - - to overwrite? - écraser ? + + PKG and Game versions match: + Les versions PKG et jeu correspondent : - - Patch detected!\nPKG Version %1 is older - Patch détecté !\nVersion PKG %1 est plus ancienne + + Would you like to overwrite? + Souhaitez-vous remplacer ? - - than installed version!: %2\nWould you like - que la version installée ! : %2\nSouhaitez-vous + + PKG Version %1 is older than installed version: + La version PKG %1 est plus ancienne que la version installée : - - to overwrite? - écraser ? + + Game is installed: + Jeu installé : - - Patch detected!\nGame is installed: %1\nWould you like - Patch détecté !\nJeu est installé : %1\nSouhaitez-vous + + Would you like to install Patch: + Souhaitez-vous installer le patch : - - to install Patch: %2? - installer le patch : %2 ? + + DLC Installation + Installation du DLC - - Game already installed\n%1\nWould you like to overwrite? - Jeu déjà installé\n%1\nSouhaitez-vous écraser ? + + Would you like to install DLC: %1? + Souhaitez-vous installer le DLC : %1 ? + + + + DLC already installed: + DLC déjà installé : + + + + Game already installed + Jeu déjà installé diff --git a/src/qt_gui/translations/hu_HU.ts b/src/qt_gui/translations/hu_HU.ts index e531df45..e5fb25a5 100644 --- a/src/qt_gui/translations/hu_HU.ts +++ b/src/qt_gui/translations/hu_HU.ts @@ -566,44 +566,54 @@ PKG kicsomagolás - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Javítás észlelve!\nA PKG és a játék verziók egyeznek: %1\nSzeretnéd - - - - to overwrite? - felülírni? - - - - Patch detected!\nPKG Version %1 is older - Javítás észlelve!\nA PKG verzió %1 régebbi - - - - than installed version!: %2\nWould you like - mint a telepített verzió: %2\nSzeretnéd - - - - to overwrite? - felülírni? + + Patch detected! + Frissítés észlelve! - Patch detected!\nGame is installed: %1\nWould you like - Javítás észlelve!\nA játék telepítve van: %1\nSzeretnéd + PKG and Game versions match: + A PKG és a játék verziói egyeznek: - to install Patch: %2? - a javítást telepíteni: %2? + Would you like to overwrite? + Szeretné felülírni? - - Game already installed\n%1\nWould you like to overwrite? - A játék már telepítve van\n%1\nSzeretnéd felülírni? + + PKG Version %1 is older than installed version: + A %1-es PKG verzió régebbi, mint a telepített verzió: + + + + Game is installed: + A játék telepítve van: + + + + Would you like to install Patch: + Szeretné telepíteni a frissítést: + + + + DLC Installation + DLC Telepítés + + + + Would you like to install DLC: %1? + Szeretné telepíteni a DLC-t: %1? + + + + DLC already installed: + DLC már telepítve: + + + + Game already installed + A játék már telepítve van diff --git a/src/qt_gui/translations/id.ts b/src/qt_gui/translations/id.ts index dde59046..b8ce27cd 100644 --- a/src/qt_gui/translations/id.ts +++ b/src/qt_gui/translations/id.ts @@ -566,44 +566,54 @@ Ekstraksi PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch terdeteksi!\nVersi PKG dan Game cocok!: %1\nApakah Anda ingin - - - - to overwrite? - menimpa? - - - - Patch detected!\nPKG Version %1 is older - Patch terdeteksi!\nVersi PKG %1 lebih lama - - - - than installed version!: %2\nWould you like - daripada versi yang terpasang!: %2\nApakah Anda ingin - - - - to overwrite? - menimpa? + + Patch detected! + Patch terdeteksi! - Patch detected!\nGame is installed: %1\nWould you like - Patch terdeteksi!\nGame terpasang: %1\nApakah Anda ingin + PKG and Game versions match: + Versi PKG dan Game cocok: - to install Patch: %2? - memasang Patch: %2? + Would you like to overwrite? + Apakah Anda ingin menimpa? - - Game already installed\n%1\nWould you like to overwrite? - Game sudah terpasang\n%1\nApakah Anda ingin menimpa? + + PKG Version %1 is older than installed version: + Versi PKG %1 lebih lama dari versi yang terpasang: + + + + Game is installed: + Game telah terpasang: + + + + Would you like to install Patch: + Apakah Anda ingin menginstal patch: + + + + DLC Installation + Instalasi DLC + + + + Would you like to install DLC: %1? + Apakah Anda ingin menginstal DLC: %1? + + + + DLC already installed: + DLC sudah terpasang: + + + + Game already installed + Game sudah terpasang diff --git a/src/qt_gui/translations/it.ts b/src/qt_gui/translations/it.ts index 4cb050b7..380a8e43 100644 --- a/src/qt_gui/translations/it.ts +++ b/src/qt_gui/translations/it.ts @@ -181,7 +181,7 @@ Install application from a .pkg file - Installa applicazione da un file .pkg file + Installa applicazione da un file .pkg @@ -236,7 +236,7 @@ List View - Visualizzazione lista + Visualizzazione Lista @@ -341,7 +341,7 @@ toolBar - barra strumenti + Barra strumenti @@ -356,8 +356,8 @@ TrophyViewer - Visualizzatore Trofei - Trophy Viewer + Trophy Viewer + Visualizzatore Trofei @@ -503,42 +503,42 @@ * Unsupported Vulkan Version - * Versi Vulkan Tidak Didukung + * Versione Vulkan non supportata Download Cheats For All Installed Games - Unduh Cheat Untuk Semua Game yang Terinstal + Scarica Trucchi per tutti i giochi installati Download Patches For All Games - Unduh Patch Untuk Semua Game + Scarica Patch per tutti i giochi Download Complete - Unduhan Selesai + Scaricamento completato You have downloaded cheats for all the games you have installed. - Anda telah mengunduh cheat untuk semua game yang telah Anda instal. + Hai scaricato trucchi per tutti i giochi installati. Patches Downloaded Successfully! - Patch Berhasil Diunduh! + Patch scaricate con successo! All Patches available for all games have been downloaded. - Semua patch yang tersedia untuk semua game telah diunduh. + Tutte le patch disponibili per tutti i giochi sono state scaricate. Games: - Game: + Giochi: @@ -553,87 +553,97 @@ Game Boot - Boot Game + Avvia Gioco Only one file can be selected! - Hanya satu file yang dapat dipilih! + Si può selezionare solo un file! PKG Extraction - Ekstraksi PKG + Estrazione file PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch terdeteksi!\nVersi PKG dan Game cocok!: %1\nApakah Anda ingin + + Patch detected! + Patch rilevata! - - to overwrite? - menimpa? + + PKG and Game versions match: + Le versioni di PKG e del gioco corrispondono: - - Patch detected!\nPKG Version %1 is older - Patch terdeteksi!\nVersi PKG %1 lebih lama + + Would you like to overwrite? + Vuoi sovrascrivere? - - than installed version!: %2\nWould you like - daripada versi yang terinstal!: %2\nApakah Anda ingin + + PKG Version %1 is older than installed version: + La versione PKG %1 è più vecchia rispetto alla versione installata: - - to overwrite? - menimpa? + + Game is installed: + Gioco installato: - - Patch detected!\nGame is installed: %1\nWould you like - Patch terdeteksi!\nGame terinstal: %1\nApakah Anda ingin + + Would you like to install Patch: + Vuoi installare la patch: - - to install Patch: %2? - menginstal Patch: %2? + + DLC Installation + Installazione DLC - - Game already installed\n%1\nWould you like to overwrite? - Game sudah terinstal\n%1\nApakah Anda ingin menimpa? + + Would you like to install DLC: %1? + Vuoi installare il DLC: %1? + + + + DLC already installed: + DLC già installato: + + + + Game already installed + Gioco già installato PKG is a patch, please install the game first! - PKG adalah patch, silakan instal game terlebih dahulu! + Questo file PKG contiene una patch. Per favore, installa prima il gioco! PKG ERROR - ERROR PKG + ERRORE PKG Extracting PKG %1/%2 - Estrazione PKG %1/%2 + Estrazione file PKG %1/%2 Extraction Finished - Ekstraksi Selesai + Estrazione Completata Game successfully installed at %1 - Game berhasil diinstal di %1 + Gioco installato correttamente in %1 File doesn't appear to be a valid PKG file - File tidak tampak sebagai file PKG yang valid + Il file sembra non essere un file PKG valido @@ -641,87 +651,87 @@ Cheats / Patches - Cheat / Patch + Trucchi / Patch defaultTextEdit_MSG - 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 + 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 No Image Available - Tidak Ada Gambar + Nessuna immagine disponibile Serial: - Serial: + Seriale: Version: - Versi: + Versione: Size: - Ukuran: + Dimensione: Select Cheat File: - Pilih File Cheat: + Seleziona File Trucchi: Repository: - Repositori: + Archivio: Download Cheats - Unduh Cheat + Scarica trucchi Delete File - Hapus File + Cancella File No files selected. - Tidak ada file yang dipilih. + Nessun file selezionato. You can delete the cheats you don't want after downloading them. - Anda dapat menghapus cheat yang tidak diinginkan setelah mengunduhnya. + Puoi cancellare i trucchi che non vuoi utilizzare dopo averli scaricati. Do you want to delete the selected file?\n%1 - Apakah Anda ingin menghapus file yang dipilih?\n%1 + Vuoi cancellare il file selezionato?\n%1 Select Patch File: - Pilih File Patch: + Seleziona File Patch: Download Patches - Unduh Patch + Scarica Patch Save - Simpan + Salva Cheats - Cheat + Trucchi @@ -731,162 +741,162 @@ Error - Kesalahan + Errore No patch selected. - Tidak ada patch yang dipilih. + Nessuna patch selezionata. Unable to open files.json for reading. - Gagal membuka files.json untuk dibaca. + Impossibile aprire il file .json per la lettura. No patch file found for the current serial. - Tidak ada file patch ditemukan untuk serial saat ini. + Nessun file patch trovato per il seriale selezionato. Unable to open the file for reading. - Gagal membuka file untuk dibaca. + Impossibile aprire il file per la lettura. Unable to open the file for writing. - Gagal membuka file untuk ditulis. + Impossibile aprire il file per la scrittura. Failed to parse XML: - Gagal mengurai XML: + Analisi XML fallita: Success - Berhasil + Successo Options saved successfully. - Opsi berhasil disimpan. + Opzioni salvate con successo. Invalid Source - Sumber Tidak Valid + Fonte non valida The selected source is invalid. - Sumber yang dipilih tidak valid. + La fonte selezionata non è valida. File Exists - File Ada + Il file è presente File already exists. Do you want to replace it? - File sudah ada. Apakah Anda ingin menggantinya? + Il file è già presente. Vuoi sostituirlo? Failed to save file: - Gagal menyimpan file: + Salvataggio file fallito: Failed to download file: - Gagal mengunduh file: + Scaricamento file fallito: Cheats Not Found - Cheat Tidak Ditemukan + Trucchi non trovati CheatsNotFound_MSG - Cheat tidak ditemukan untuk game ini dalam versi repositori yang dipilih, coba repositori lain atau versi game yang berbeda. + Non sono stati trovati trucchi per questa versione del gioco nell'archivio selezionato, prova un altro archivio o una versione diversa del gioco. Cheats Downloaded Successfully - Cheat Berhasil Diunduh + Trucchi scaricati con successo! CheatsDownloadedSuccessfully_MSG - 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. + 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. Failed to save: - Gagal menyimpan: + Salvataggio fallito: Failed to download: - Gagal mengunduh: + Impossibile scaricare: Download Complete - Unduhan Selesai + Scaricamento completo DownloadComplete_MSG - 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. + 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. Failed to parse JSON data from HTML. - Gagal mengurai data JSON dari HTML. + Impossibile analizzare i dati JSON dall'HTML. Failed to retrieve HTML page. - Gagal mengambil halaman HTML. + Impossibile recuperare la pagina HTML. Failed to open file: - Gagal membuka file: + Impossibile aprire file: XML ERROR: - KESALAHAN XML: + ERRORE XML: Failed to open files.json for writing - Gagal membuka files.json untuk menulis + Impossibile aprire i file .json per la scrittura Author: - Penulis: + Autore: Directory does not exist: - Direktori tidak ada: + La cartella non esiste: Failed to open files.json for reading. - Gagal membuka files.json untuk dibaca. + Impossibile aprire i file .json per la lettura. Name: - Nama: + Nome: - \ No newline at end of file + diff --git a/src/qt_gui/translations/ja_JP.ts b/src/qt_gui/translations/ja_JP.ts index c3eaaba1..3d62de0d 100644 --- a/src/qt_gui/translations/ja_JP.ts +++ b/src/qt_gui/translations/ja_JP.ts @@ -566,44 +566,54 @@ PKG抽出 - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - パッチが検出されました！\nPKGとゲームバージョンが一致しています！: %1\n上書きしますか - - - - to overwrite? - 上書きしますか？ - - - - Patch detected!\nPKG Version %1 is older - パッチが検出されました！\nPKGバージョン %1 は古い - - - - than installed version!: %2\nWould you like - インストールされているバージョンよりも古いです！: %2\n上書きしますか - - - - to overwrite? - 上書きしますか？ + + Patch detected! + パッチが検出されました！ - Patch detected!\nGame is installed: %1\nWould you like - パッチが検出されました！\nゲームがインストールされています: %1\nインストールしますか + PKG and Game versions match: + PKGとゲームのバージョンが一致しています: - to install Patch: %2? - パッチをインストールしますか: %2? + Would you like to overwrite? + 上書きしてもよろしいですか？ - - Game already installed\n%1\nWould you like to overwrite? - ゲームはすでにインストールされています\n%1\n上書きしますか？ + + PKG Version %1 is older than installed version: + PKGバージョン %1 はインストールされているバージョンよりも古いです: + + + + Game is installed: + ゲームはインストール済みです: + + + + Would you like to install Patch: + パッチをインストールしてもよろしいですか: + + + + DLC Installation + DLCのインストール + + + + Would you like to install DLC: %1? + DLCをインストールしてもよろしいですか: %1? + + + + DLC already installed: + DLCはすでにインストールされています: + + + + Game already installed + ゲームはすでにインストールされています diff --git a/src/qt_gui/translations/ko_KR.ts b/src/qt_gui/translations/ko_KR.ts index 579c6ca7..f7f171dc 100644 --- a/src/qt_gui/translations/ko_KR.ts +++ b/src/qt_gui/translations/ko_KR.ts @@ -566,44 +566,54 @@ PKG Extraction - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch detected!\nPKG and Game versions match!: %1\nWould you like - - - - to overwrite? - to overwrite? - - - - Patch detected!\nPKG Version %1 is older - Patch detected!\nPKG Version %1 is older - - - - than installed version!: %2\nWould you like - than installed version!: %2\nWould you like - - - - to overwrite? - to overwrite? + + Patch detected! + Patch detected! - Patch detected!\nGame is installed: %1\nWould you like - Patch detected!\nGame is installed: %1\nWould you like + PKG and Game versions match: + PKG and Game versions match: - to install Patch: %2? - to install Patch: %2? + Would you like to overwrite? + Would you like to overwrite? - - Game already installed\n%1\nWould you like to overwrite? - Game already installed\n%1\nWould you like to overwrite? + + PKG Version %1 is older than installed version: + PKG Version %1 is older than installed version: + + + + Game is installed: + Game is installed: + + + + Would you like to install Patch: + Would you like to install Patch: + + + + DLC Installation + DLC Installation + + + + Would you like to install DLC: %1? + Would you like to install DLC: %1? + + + + DLC already installed: + DLC already installed: + + + + Game already installed + Game already installed diff --git a/src/qt_gui/translations/lt_LT.ts b/src/qt_gui/translations/lt_LT.ts index c94edb74..7aa4402e 100644 --- a/src/qt_gui/translations/lt_LT.ts +++ b/src/qt_gui/translations/lt_LT.ts @@ -566,44 +566,54 @@ PKG ištraukimas - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Pataisa aptikta!\nPKG ir žaidimo versijos atitinka!: %1\nAr norėtumėte - - - - to overwrite? - perrašyti? - - - - Patch detected!\nPKG Version %1 is older - Pataisa aptikta!\nPKG versija %1 yra senesnė - - - - than installed version!: %2\nWould you like - nei įdiegta versija!: %2\nAr norėtumėte - - - - to overwrite? - perrašyti? + + Patch detected! + Rasta atnaujinimą! - Patch detected!\nGame is installed: %1\nWould you like - Pataisa aptikta!\nŽaidimas įdiegtas: %1\nAr norėtumėte + PKG and Game versions match: + PKG ir žaidimo versijos sutampa: - to install Patch: %2? - įdiegti pataisą: %2? + Would you like to overwrite? + Ar norite perrašyti? - - Game already installed\n%1\nWould you like to overwrite? - Žaidimas jau įdiegtas\n%1\nAr norėtumėte perrašyti? + + PKG Version %1 is older than installed version: + PKG versija %1 yra senesnė nei įdiegta versija: + + + + Game is installed: + Žaidimas įdiegtas: + + + + Would you like to install Patch: + Ar norite įdiegti atnaujinimą: + + + + DLC Installation + DLC diegimas + + + + Would you like to install DLC: %1? + Ar norite įdiegti DLC: %1? + + + + DLC already installed: + DLC jau įdiegtas: + + + + Game already installed + Žaidimas jau įdiegtas diff --git a/src/qt_gui/translations/nb.ts b/src/qt_gui/translations/nb.ts index 3c5401a2..76cad45b 100644 --- a/src/qt_gui/translations/nb.ts +++ b/src/qt_gui/translations/nb.ts @@ -566,44 +566,54 @@ PKG-ekstraksjon - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Oppdatering oppdaget!\nPKG og spillversjoner stemmer!: %1\nØnsker du å + + Patch detected! + Oppdatering oppdaget! - - to overwrite? - overskrive? + + PKG and Game versions match: + PKG- og spillversjoner stemmer overens: - - Patch detected!\nPKG Version %1 is older - Oppdatering oppdaget!\nPKG-versjon %1 er eldre + + Would you like to overwrite? + Ønsker du å overskrive? - - than installed version!: %2\nWould you like - enn installert versjon!: %2\nØnsker du å + + PKG Version %1 is older than installed version: + PKG-versjon %1 er eldre enn installert versjon: - - to overwrite? - overskrive? + + Game is installed: + Spillet er installert: - - Patch detected!\nGame is installed: %1\nWould you like - Oppdatering oppdaget!\nSpillet er installert: %1\nØnsker du å + + Would you like to install Patch: + Ønsker du å installere oppdateringen: - - to install Patch: %2? - installere oppdateringen: %2? + + DLC Installation + DLC-installasjon - - Game already installed\n%1\nWould you like to overwrite? - Spill allerede installert\n%1\nØnsker du å overskrive? + + Would you like to install DLC: %1? + Ønsker du å installere DLC: %1? + + + + DLC already installed: + DLC allerede installert: + + + + Game already installed + Spillet er allerede installert diff --git a/src/qt_gui/translations/nl.ts b/src/qt_gui/translations/nl.ts index 8b55b0e2..b0046047 100644 --- a/src/qt_gui/translations/nl.ts +++ b/src/qt_gui/translations/nl.ts @@ -566,44 +566,54 @@ PKG-extractie - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch gedetecteerd!\nPKG en spelversies komen overeen!: %1\nWil je + + Patch detected! + Patch gedetecteerd! - - to overwrite? - overschrijven? + + PKG and Game versions match: + PKG- en gameversies komen overeen: - - Patch detected!\nPKG Version %1 is older - Patch gedetecteerd!\nPKG-versie %1 is ouder + + Would you like to overwrite? + Wilt u overschrijven? - - than installed version!: %2\nWould you like - dan de geïnstalleerde versie!: %2\nWil je + + PKG Version %1 is older than installed version: + PKG-versie %1 is ouder dan de geïnstalleerde versie: - - to overwrite? - overschrijven? + + Game is installed: + Game is geïnstalleerd: - - Patch detected!\nGame is installed: %1\nWould you like - Patch gedetecteerd!\nSpel is geïnstalleerd: %1\nWil je + + Would you like to install Patch: + Wilt u de patch installeren: - - to install Patch: %2? - de patch installeren: %2? + + DLC Installation + DLC-installatie - - Game already installed\n%1\nWould you like to overwrite? - Spel al geïnstalleerd\n%1\nWil je het overschrijven? + + Would you like to install DLC: %1? + Wilt u DLC installeren: %1? + + + + DLC already installed: + DLC al geïnstalleerd: + + + + Game already installed + Game al geïnstalleerd diff --git a/src/qt_gui/translations/pl_PL.ts b/src/qt_gui/translations/pl_PL.ts index 41c06f89..80f0aa5c 100644 --- a/src/qt_gui/translations/pl_PL.ts +++ b/src/qt_gui/translations/pl_PL.ts @@ -566,44 +566,54 @@ Ekstrakcja PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Wykryto poprawkę!\nWersje PKG i gry pasują do siebie!: %1\nCzy chcesz + + Patch detected! + Wykryto łatkę! - - to overwrite? - nadpisać? + + PKG and Game versions match: + Wersje PKG i gry są zgodne: - - Patch detected!\nPKG Version %1 is older - Wykryto poprawkę!\nWersja PKG %1 jest starsza + + Would you like to overwrite? + Czy chcesz nadpisać? - - than installed version!: %2\nWould you like - niż zainstalowana wersja!: %2\nCzy chcesz + + PKG Version %1 is older than installed version: + Wersja PKG %1 jest starsza niż zainstalowana wersja: - - to overwrite? - nadpisać? + + Game is installed: + Gra jest zainstalowana: - - Patch detected!\nGame is installed: %1\nWould you like - Wykryto poprawkę!\nGra jest zainstalowana: %1\nCzy chcesz + + Would you like to install Patch: + Czy chcesz zainstalować łatkę: - - to install Patch: %2? - zainstalować poprawkę: %2? + + DLC Installation + Instalacja DLC - - Game already installed\n%1\nWould you like to overwrite? - Gra już zainstalowana\n%1\nCzy chcesz ją nadpisać? + + Would you like to install DLC: %1? + Czy chcesz zainstalować DLC: %1? + + + + DLC already installed: + DLC już zainstalowane: + + + + Game already installed + Gra już zainstalowana diff --git a/src/qt_gui/translations/pt_BR.ts b/src/qt_gui/translations/pt_BR.ts index c198d1fd..8b4538b9 100644 --- a/src/qt_gui/translations/pt_BR.ts +++ b/src/qt_gui/translations/pt_BR.ts @@ -37,7 +37,7 @@ Loading game list, please wait :3 - Carregando lista de jogos, por favor aguarde :3 + Carregando a lista de jogos, por favor aguarde :3 @@ -256,7 +256,7 @@ Download Cheats/Patches - Baixar Trapaças / Patches + Baixar Cheats/Patches @@ -291,7 +291,7 @@ Game List Mode - Modo de Lista de Jogos + Modo da Lista de Jogos @@ -425,7 +425,7 @@ Log Filter - Filtro + Filtro do Registro @@ -508,12 +508,12 @@ Download Cheats For All Installed Games - Baixar Trapaças para todos os jogos instalados + Baixar Cheats para Todos os Jogos Instalados Download Patches For All Games - Baixar Patches para todos os jogos + Baixar Patches para Todos os Jogos @@ -523,7 +523,7 @@ You have downloaded cheats for all the games you have installed. - Você baixou trapaças para todos os jogos que instalou. + Você baixou cheats para todos os jogos que instalou. @@ -566,54 +566,64 @@ Extração de PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch detectado!\nVersões PKG e do Jogo correspondem!: %1\nGostaria de - - - - to overwrite? - substituir? - - - - Patch detected!\nPKG Version %1 is older - Patch detectado!\nVersão PKG %1 é mais antiga - - - - than installed version!: %2\nWould you like - do que a versão instalada!: %2\nGostaria de - - - - to overwrite? - substituir? + + Patch detected! + Atualização detectada! - Patch detected!\nGame is installed: %1\nWould you like - Patch detectado!\nJogo está instalado: %1\nGostaria de + PKG and Game versions match: + As versões do PKG e do Jogo são igual: - to install Patch: %2? - instalar o Patch: %2? + Would you like to overwrite? + Gostaria de substituir? - - Game already installed\n%1\nWould you like to overwrite? - Jogo já instalado\n%1\nGostaria de substituir? + + PKG Version %1 is older than installed version: + Versão do PKG %1 é mais antiga do que a versão instalada: + + + + Game is installed: + Jogo instalado: + + + + Would you like to install Patch: + Você gostaria de instalar a atualização: + + + + DLC Installation + Instalação de DLC + + + + Would you like to install DLC: %1? + Você gostaria de instalar o DLC: %1? + + + + DLC already installed: + DLC já instalada: + + + + Game already installed + O jogo já está instalado: PKG is a patch, please install the game first! - PKG é um patch, por favor, instale o jogo primeiro! + O PKG é um patch, por favor, instale o jogo primeiro! PKG ERROR - ERRO PKG + ERRO de PKG @@ -641,12 +651,12 @@ Cheats / Patches - Trapaças / Patches + Cheats / Patches defaultTextEdit_MSG - 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 + 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 @@ -656,7 +666,7 @@ Serial: - Série: + Serial: @@ -671,7 +681,7 @@ Select Cheat File: - Selecione o Arquivo de Trapaça: + Selecione o Arquivo de Cheat: @@ -681,7 +691,7 @@ Download Cheats - Baixar Trapaças + Baixar Cheats @@ -696,7 +706,7 @@ You can delete the cheats you don't want after downloading them. - Você pode excluir as trapaças que não deseja após baixá-las. + Você pode excluir os cheats que não deseja após baixá-las. @@ -721,7 +731,7 @@ Cheats - Trapaças + Cheats @@ -746,7 +756,7 @@ No patch file found for the current serial. - Nenhum arquivo de patch encontrado para a série atual. + Nenhum arquivo de patch encontrado para o serial atual. @@ -806,22 +816,22 @@ Cheats Not Found - Trapaças Não Encontradas + Cheats Não Encontrados CheatsNotFound_MSG - 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. + Nenhum cheat encontrado para este jogo nesta versão do repositório selecionado, tente outro repositório ou uma versão diferente do jogo. Cheats Downloaded Successfully - Trapaças Baixadas com Sucesso + Cheats Baixados com Sucesso CheatsDownloadedSuccessfully_MSG - 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. + 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. @@ -841,7 +851,7 @@ DownloadComplete_MSG - 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. + 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. @@ -861,7 +871,7 @@ XML ERROR: - ERRO XML: + ERRO de XML: @@ -876,7 +886,7 @@ Directory does not exist: - Diretório não existe: + O Diretório não existe: diff --git a/src/qt_gui/translations/ro_RO.ts b/src/qt_gui/translations/ro_RO.ts index 3463182f..8b2fda0c 100644 --- a/src/qt_gui/translations/ro_RO.ts +++ b/src/qt_gui/translations/ro_RO.ts @@ -566,44 +566,54 @@ Extracție PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Patch detectat!\nVersiunile PKG și Joc se potrivesc!: %1\nAi dori să + + Patch detected! + Patch detectat! - - to overwrite? - să suprascrii? + + PKG and Game versions match: + Versiunile PKG și ale jocului sunt compatibile: - - Patch detected!\nPKG Version %1 is older - Patch detectat!\nVersiunea PKG %1 este mai veche + + Would you like to overwrite? + Doriți să suprascrieți? - - than installed version!: %2\nWould you like - decât versiunea instalată!: %2\nAi dori să + + PKG Version %1 is older than installed version: + Versiunea PKG %1 este mai veche decât versiunea instalată: - - to overwrite? - să suprascrii? + + Game is installed: + Jocul este instalat: - - Patch detected!\nGame is installed: %1\nWould you like - Patch detectat!\nJocul este instalat: %1\nAi dori să + + Would you like to install Patch: + Doriți să instalați patch-ul: - - to install Patch: %2? - să instalezi Patch-ul: %2? + + DLC Installation + Instalare DLC - - Game already installed\n%1\nWould you like to overwrite? - Jocul este deja instalat\n%1\nAi dori să suprascrii? + + Would you like to install DLC: %1? + Doriți să instalați DLC-ul: %1? + + + + DLC already installed: + DLC deja instalat: + + + + Game already installed + Jocul deja instalat diff --git a/src/qt_gui/translations/ru_RU.ts b/src/qt_gui/translations/ru_RU.ts index a71533a3..9e3446ad 100644 --- a/src/qt_gui/translations/ru_RU.ts +++ b/src/qt_gui/translations/ru_RU.ts @@ -93,7 +93,7 @@ Cheats / Patches - Читы / Патчи + Читы и патчи @@ -256,7 +256,7 @@ Download Cheats/Patches - Скачать Читы / Патчи + Скачать читы или патчи @@ -566,45 +566,55 @@ Извлечение PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Обнаружен патч!\nВерсии PKG и игры совпадают!: %1\nХотите - - - - to overwrite? - перезаписать? - - - - Patch detected!\nPKG Version %1 is older - Обнаружен патч!\nВерсия PKG %1 устарела - - - - than installed version!: %2\nWould you like - по сравнению с установленной версией!: %2\nХотите - - - - to overwrite? - перезаписать? + + Patch detected! + Обнаружен патч! - Patch detected!\nGame is installed: %1\nWould you like - Обнаружен патч!\nИгра установлена: %1\nХотите + PKG and Game versions match: + Версии PKG и игры совпадают: - to install Patch: %2? - установить патч: %2? + Would you like to overwrite? + Хотите перезаписать? - - Game already installed\n%1\nWould you like to overwrite? - Игра уже установлена\n%1\nХотите перезаписать? + + PKG Version %1 is older than installed version: + Версия PKG %1 старее установленной версии: + + + Game is installed: + Игра установлена: + + + + Would you like to install Patch: + Хотите установить патч: + + + + DLC Installation + Установка DLC + + + + Would you like to install DLC: %1? + Вы хотите установить DLC: %1?? + + + + DLC already installed: + DLC уже установлен: + + + + Game already installed + Игра уже установлена + PKG is a patch, please install the game first! @@ -628,7 +638,7 @@ Game successfully installed at %1 - Игра успешно установлена по адресу %1 + Игра успешно установлена в %1 @@ -641,13 +651,13 @@ Cheats / Patches - Читы / Патчи + Читы и патчи defaultTextEdit_MSG - 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 să le utilizați și salvați selecția.\n\nDeoarece nu dezvoltăm Cheats/Patches,\nte rugăm să raportezi problemele autorului cheat-ului.\n\nAi creat un nou cheat? Vizitează:\nhttps://github.com/shadps4-emu/ps4_cheats - + Читы и патчи экспериментальны.\nИспользуйте с осторожностью.\n\nСкачивайте читы, выбрав репозиторий и нажав на кнопку загрузки.\nВо вкладке "Патчи" вы можете скачать все патчи сразу, выбирать какие вы хотите использовать, и сохранять выбор.\n\nПоскольку мы не разрабатываем читы/патчи,\nпожалуйста сообщайте о проблемах автору чита/патча.\n\nСоздали новый чит? Посетите:\nhttps://github.com/shadps4-emu/ps4_cheats + No Image Available diff --git a/src/qt_gui/translations/tr_TR.ts b/src/qt_gui/translations/tr_TR.ts index 514b9af7..e11a2d96 100644 --- a/src/qt_gui/translations/tr_TR.ts +++ b/src/qt_gui/translations/tr_TR.ts @@ -566,44 +566,54 @@ PKG Çıkartma - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Yama tespit edildi!\nPKG ve Oyun sürümleri uyuyor!: %1\nÜzerine yazmak ister misiniz? - - - - to overwrite? - üzerine yazmak? - - - - Patch detected!\nPKG Version %1 is older - Yama tespit edildi!\nPKG Sürümü %1 daha eski - - - - than installed version!: %2\nWould you like - yüklü sürümden!: %2\nÜzerine yazmak ister misiniz? - - - - to overwrite? - üzerine yazmak? + + Patch detected! + Yamanın tespit edildi! - Patch detected!\nGame is installed: %1\nWould you like - Yama tespit edildi!\nOyun yüklü: %1\nÜzerine yazmak ister misiniz? + PKG and Game versions match: + PKG ve oyun sürümleri uyumlu: - to install Patch: %2? - Yamayı kurmak ister misiniz: %2? + Would you like to overwrite? + Üzerine yazmak ister misiniz? - - Game already installed\n%1\nWould you like to overwrite? - Oyun zaten yüklü\n%1\nÜzerine yazmak ister misiniz? + + PKG Version %1 is older than installed version: + PKG Sürümü %1, kurulu sürümden daha eski: + + + + Game is installed: + Oyun yüklendi: + + + + Would you like to install Patch: + Yamanın yüklenmesini ister misiniz: + + + + DLC Installation + DLC Yükleme + + + + Would you like to install DLC: %1? + DLC'yi yüklemek ister misiniz: %1? + + + + DLC already installed: + DLC zaten yüklü: + + + + Game already installed + Oyun zaten yüklü diff --git a/src/qt_gui/translations/vi_VN.ts b/src/qt_gui/translations/vi_VN.ts index 977b6760..aead45a6 100644 --- a/src/qt_gui/translations/vi_VN.ts +++ b/src/qt_gui/translations/vi_VN.ts @@ -566,45 +566,55 @@ Giải nén PKG - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - Đã phát hiện bản vá!\nPhiên bản PKG và trò chơi khớp!: %1\nBạn có muốn + + Patch detected! + Đã phát hiện bản vá! - - to overwrite? - ghi đè không? + + PKG and Game versions match: + Các phiên bản PKG và trò chơi khớp nhau: - - Patch detected!\nPKG Version %1 is older - Đã phát hiện bản vá!\nPhiên bản PKG %1 cũ hơn + + Would you like to overwrite? + Bạn có muốn ghi đè không? - - than installed version!: %2\nWould you like - so với phiên bản đã cài đặt!: %2\nBạn có muốn + + PKG Version %1 is older than installed version: + Phiên bản PKG %1 cũ hơn phiên bản đã cài đặt: - - to overwrite? - ghi đè không? + + Game is installed: + Trò chơi đã được cài đặt: - - Patch detected!\nGame is installed: %1\nWould you like - Đã phát hiện bản vá!\nTrò chơi đã được cài đặt: %1\nBạn có muốn + + Would you like to install Patch: + Bạn có muốn cài đặt bản vá: - - to install Patch: %2? - cài đặt bản vá: %2? + + DLC Installation + Cài đặt DLC - - Game already installed\n%1\nWould you like to overwrite? - Trò chơi đã được cài đặt\n%1\nBạn có muốn ghi đè không? + + Would you like to install DLC: %1? + Bạn có muốn cài đặt DLC: %1? + + + DLC already installed: + DLC đã được cài đặt: + + + + Game already installed + Trò chơi đã được cài đặt + PKG is a patch, please install the game first! diff --git a/src/qt_gui/translations/zh_CN.ts b/src/qt_gui/translations/zh_CN.ts index d5ace320..a1b2523b 100644 --- a/src/qt_gui/translations/zh_CN.ts +++ b/src/qt_gui/translations/zh_CN.ts @@ -566,44 +566,54 @@ PKG 解压 - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - 检测到补丁！\nPKG 和游戏版本匹配：%1\n您想要 + + Patch detected! + 检测到补丁！ - - to overwrite? - 覆盖吗？ + + PKG and Game versions match: + PKG 和游戏版本匹配: - - Patch detected!\nPKG Version %1 is older - 检测到补丁！\nPKG 版本 %1 较旧 + + Would you like to overwrite? + 您想要覆盖吗？ - - than installed version!: %2\nWould you like - 与已安装版本相比：%2\n您想要 + + PKG Version %1 is older than installed version: + PKG 版本 %1 比已安装版本更旧: - - to overwrite? - 覆盖吗？ + + Game is installed: + 游戏已安装: - - Patch detected!\nGame is installed: %1\nWould you like - 检测到补丁！\n游戏已安装：%1\n您想要 + + Would you like to install Patch: + 您想安装补丁吗: - - to install Patch: %2? - 安装补丁：%2？ + + DLC Installation + DLC 安装 - - Game already installed\n%1\nWould you like to overwrite? - 游戏已安装\n%1\n您想要覆盖吗？ + + Would you like to install DLC: %1? + 您想安装 DLC: %1 吗？ + + + + DLC already installed: + DLC 已经安装: + + + + Game already installed + 游戏已经安装 diff --git a/src/qt_gui/translations/zh_TW.ts b/src/qt_gui/translations/zh_TW.ts index 4a9da9f3..3836ed18 100644 --- a/src/qt_gui/translations/zh_TW.ts +++ b/src/qt_gui/translations/zh_TW.ts @@ -566,44 +566,54 @@ PKG 解壓縮 - - Patch detected!\nPKG and Game versions match!: %1\nWould you like - 偵測到修補檔！\nPKG 和遊戲版本匹配！: %1\n您是否希望 + + Patch detected! + 檢測到補丁！ - - to overwrite? - 覆蓋嗎？ + + PKG and Game versions match: + PKG 和遊戲版本匹配: - - Patch detected!\nPKG Version %1 is older - 偵測到修補檔！\nPKG 版本 %1 較舊 + + Would you like to overwrite? + 您想要覆蓋嗎？ - - than installed version!: %2\nWould you like - 比安裝的版本舊！: %2\n您是否希望 + + PKG Version %1 is older than installed version: + PKG 版本 %1 比已安裝版本更舊: - - to overwrite? - 覆蓋嗎？ + + Game is installed: + 遊戲已安裝: - - Patch detected!\nGame is installed: %1\nWould you like - 偵測到修補檔！\n遊戲已安裝: %1\n您是否希望 + + Would you like to install Patch: + 您想要安裝補丁嗎: - - to install Patch: %2? - 安裝修補檔: %2？ + + DLC Installation + DLC 安裝 - - Game already installed\n%1\nWould you like to overwrite? - 遊戲已經安裝\n%1\n您是否希望覆蓋？ + + Would you like to install DLC: %1? + 您想要安裝 DLC: %1 嗎？ + + + + DLC already installed: + DLC 已經安裝: + + + + Game already installed + 遊戲已經安裝 diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index 09a9fd62..98eac081 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -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 EmitSPIRV(const Profile& profile, IR::Program& program, u32& binding) { - EmitContext ctx{profile, program, binding}; +std::vector 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) { diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.h b/src/shader_recompiler/backend/spirv/emit_spirv.h index e513975b..4c862185 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.h +++ b/src/shader_recompiler/backend/spirv/emit_spirv.h @@ -9,7 +9,7 @@ namespace Shader::Backend::SPIRV { -[[nodiscard]] std::vector EmitSPIRV(const Profile& profile, IR::Program& program, +[[nodiscard]] std::vector EmitSPIRV(const Profile& profile, const IR::Program& program, u32& binding); } // namespace Shader::Backend::SPIRV diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp index 03fc52ff..7bdc98de 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp @@ -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 -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 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 @@ -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 -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 diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index 994c2847..51315139 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -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 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()) { - 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); - 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); + 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 = 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(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; } - - image_u32 = TypePointer(spv::StorageClass::Image, U32[1]); - + 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]); diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h index 5a09c411..d3646382 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -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 buffers; + boost::container::small_vector texture_buffers; boost::container::small_vector images; boost::container::small_vector samplers; @@ -238,6 +245,7 @@ private: void DefineOutputs(); void DefinePushDataBlock(); void DefineBuffers(); + void DefineTextureBuffers(); void DefineImagesAndSamplers(); void DefineSharedMemory(); diff --git a/src/shader_recompiler/frontend/translate/data_share.cpp b/src/shader_recompiler/frontend/translate/data_share.cpp index 7580f744..aa9b49b6 100644 --- a/src/shader_recompiler/frontend/translate/data_share.cpp +++ b/src/shader_recompiler/frontend/translate/data_share.cpp @@ -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}); } diff --git a/src/shader_recompiler/frontend/translate/export.cpp b/src/shader_recompiler/frontend/translate/export.cpp index 889de21b..d80de002 100644 --- a/src/shader_recompiler/frontend/translate/export.cpp +++ b/src/shader_recompiler/frontend/translate/export.cpp @@ -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())); } diff --git a/src/shader_recompiler/frontend/translate/scalar_alu.cpp b/src/shader_recompiler/frontend/translate/scalar_alu.cpp index da74f901..af258cd1 100644 --- a/src/shader_recompiler/frontend/translate/scalar_alu.cpp +++ b/src/shader_recompiler/frontend/translate/scalar_alu.cpp @@ -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: diff --git a/src/shader_recompiler/frontend/translate/translate.cpp b/src/shader_recompiler/frontend/translate/translate.cpp index e59cd565..eb86310b 100644 --- a/src/shader_recompiler/frontend/translate/translate.cpp +++ b/src/shader_recompiler/frontend/translate/translate.cpp @@ -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(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: diff --git a/src/shader_recompiler/frontend/translate/translate.h b/src/shader_recompiler/frontend/translate/translate.h index 8cbf7357..8d418421 100644 --- a/src/shader_recompiler/frontend/translate/translate.h +++ b/src/shader_recompiler/frontend/translate/translate.h @@ -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); diff --git a/src/shader_recompiler/frontend/translate/vector_alu.cpp b/src/shader_recompiler/frontend/translate/vector_alu.cpp index 274dcff1..13a8342d 100644 --- a/src/shader_recompiler/frontend/translate/vector_alu.cpp +++ b/src/shader_recompiler/frontend/translate/vector_alu.cpp @@ -415,14 +415,20 @@ void Translator::V_ADDC_U32(const GcnInst& inst) { const auto src0 = GetSrc(inst.src[0]); const auto src1 = GetSrc(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}; diff --git a/src/shader_recompiler/frontend/translate/vector_memory.cpp b/src/shader_recompiler/frontend/translate/vector_memory.cpp index b88cfc46..73530dad 100644 --- a/src/shader_recompiler/frontend/translate/vector_memory.cpp +++ b/src/shader_recompiler/frontend/translate/vector_memory.cpp @@ -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(mtbuf.dfmt)); - info.nfmt.Assign(static_cast(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(mtbuf.dfmt); + const auto nfmt = static_cast(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(mtbuf.dfmt)); - info.nfmt.Assign(static_cast(mtbuf.nfmt)); + const auto dmft = static_cast(mtbuf.dfmt); + const auto nfmt = static_cast(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(src_reg); + value = ir.GetVectorReg(src_reg); break; case 2: - value = ir.CompositeConstruct(ir.GetVectorReg(src_reg), - ir.GetVectorReg(src_reg + 1)); + value = ir.CompositeConstruct(ir.GetVectorReg(src_reg), + ir.GetVectorReg(src_reg + 1)); break; case 3: - value = ir.CompositeConstruct(ir.GetVectorReg(src_reg), - ir.GetVectorReg(src_reg + 1), - ir.GetVectorReg(src_reg + 2)); + value = ir.CompositeConstruct(ir.GetVectorReg(src_reg), + ir.GetVectorReg(src_reg + 1), + ir.GetVectorReg(src_reg + 2)); break; case 4: - value = ir.CompositeConstruct(ir.GetVectorReg(src_reg), - ir.GetVectorReg(src_reg + 1), - ir.GetVectorReg(src_reg + 2), - ir.GetVectorReg(src_reg + 3)); + value = ir.CompositeConstruct( + ir.GetVectorReg(src_reg), ir.GetVectorReg(src_reg + 1), + ir.GetVectorReg(src_reg + 2), ir.GetVectorReg(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); - } + ir.StoreBuffer(num_dwords, handle, address, value, info); +} + +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 comps{}; + for (u32 i = 0; i < num_dwords; i++) { + comps[i] = ir.GetVectorReg(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); } -// TODO: U64 void Translator::BUFFER_ATOMIC(AtomicOp op, const GcnInst& inst) { const auto& mubuf = inst.control.mubuf; const IR::VectorReg vaddr{inst.src[0].code}; diff --git a/src/shader_recompiler/ir/ir_emitter.cpp b/src/shader_recompiler/ir/ir_emitter.cpp index 65de98b7..473ae4f6 100644 --- a/src/shader_recompiler/ir/ir_emitter.cpp +++ b/src/shader_recompiler/ir/ir_emitter.cpp @@ -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: - 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); - } +Value IREmitter::LoadBufferFormat(const Value& handle, const Value& address, BufferInstInfo info) { + return Inst(Opcode::LoadBufferFormatF32, Flags{info}, handle, address); } 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: - 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); - } +void IREmitter::StoreBufferFormat(const Value& handle, const Value& address, const Value& data, + BufferInstInfo info) { + Inst(Opcode::StoreBufferFormatF32, Flags{info}, handle, address, data); } U32 IREmitter::LaneId() { diff --git a/src/shader_recompiler/ir/ir_emitter.h b/src/shader_recompiler/ir/ir_emitter.h index a60f4c28..de8fe450 100644 --- a/src/shader_recompiler/ir/ir_emitter.h +++ b/src/shader_recompiler/ir/ir_emitter.h @@ -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); diff --git a/src/shader_recompiler/ir/microinstruction.cpp b/src/shader_recompiler/ir/microinstruction.cpp index a8c8b073..d6ef49cf 100644 --- a/src/shader_recompiler/ir/microinstruction.cpp +++ b/src/shader_recompiler/ir/microinstruction.cpp @@ -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: diff --git a/src/shader_recompiler/ir/opcodes.inc b/src/shader_recompiler/ir/opcodes.inc index a49ea1c7..1e33d6d4 100644 --- a/src/shader_recompiler/ir/opcodes.inc +++ b/src/shader_recompiler/ir/opcodes.inc @@ -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 diff --git a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp index ace6a37d..f446ac47 100644 --- a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp +++ b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp @@ -3,6 +3,7 @@ #include #include +#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::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(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(); - 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(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(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); } diff --git a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp index 52087a65..7105f01f 100644 --- a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp +++ b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp @@ -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; } diff --git a/src/shader_recompiler/ir/program.h b/src/shader_recompiler/ir/program.h index eff933f2..f7abba64 100644 --- a/src/shader_recompiler/ir/program.h +++ b/src/shader_recompiler/ir/program.h @@ -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 ins_list; - Info info; + Info& info; }; [[nodiscard]] std::string DumpProgram(const Program& program); diff --git a/src/shader_recompiler/ir/reg.h b/src/shader_recompiler/ir/reg.h index 7868a5a3..fba04f33 100644 --- a/src/shader_recompiler/ir/reg.h +++ b/src/shader_recompiler/ir/reg.h @@ -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 { diff --git a/src/shader_recompiler/recompiler.cpp b/src/shader_recompiler/recompiler.cpp index 0efac4ff..dfcf9ed1 100644 --- a/src/shader_recompiler/recompiler.cpp +++ b/src/shader_recompiler/recompiler.cpp @@ -29,7 +29,7 @@ IR::BlockList GenerateBlocks(const IR::AbstractSyntaxList& syntax_list) { IR::Program TranslateProgram(Common::ObjectPool& inst_pool, Common::ObjectPool& block_pool, std::span 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& 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& 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()); diff --git a/src/shader_recompiler/recompiler.h b/src/shader_recompiler/recompiler.h index 34e958a1..3a229518 100644 --- a/src/shader_recompiler/recompiler.h +++ b/src/shader_recompiler/recompiler.h @@ -13,7 +13,7 @@ struct Profile; [[nodiscard]] IR::Program TranslateProgram(Common::ObjectPool& inst_pool, Common::ObjectPool& block_pool, - std::span code, const Info&& info, + std::span code, Info& info, const Profile& profile); } // namespace Shader diff --git a/src/shader_recompiler/runtime_info.h b/src/shader_recompiler/runtime_info.h index 7d36dbe1..77c57e94 100644 --- a/src/shader_recompiler/runtime_info.h +++ b/src/shader_recompiler/runtime_info.h @@ -4,6 +4,7 @@ #pragma once #include +#include #include #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; +using BufferResourceList = boost::container::small_vector; + +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; 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; +using ImageResourceList = boost::container::small_vector; 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; +using SamplerResourceList = boost::container::small_vector; 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 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(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(sgpr_base, dword_offset); +} + +constexpr AmdGpu::Image ImageResource::GetSharp(const Info& info) const noexcept { + return info.ReadUd(sgpr_base, dword_offset); +} + +constexpr AmdGpu::Sampler SamplerResource::GetSharp(const Info& info) const noexcept { return inline_sampler ? inline_sampler : info.ReadUd(sgpr_base, dword_offset); } diff --git a/src/video_core/amdgpu/liverpool.cpp b/src/video_core/amdgpu/liverpool.cpp index 8570a290..e61f8cec 100644 --- a/src/video_core/amdgpu/liverpool.cpp +++ b/src/video_core/amdgpu/liverpool.cpp @@ -383,6 +383,22 @@ Liverpool::Task Liverpool::ProcessGraphics(std::span dcb, std::span(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(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(header); regs.num_instances.num_instances = num_instances->num_instances; @@ -399,6 +415,12 @@ Liverpool::Task Liverpool::ProcessGraphics(std::span dcb, std::spannum_indices; break; } + case PM4ItOpcode::SetBase: { + const auto* set_base = reinterpret_cast(header); + ASSERT(set_base->base_index == PM4CmdSetBase::BaseIndex::DrawIndexIndirPatchTable); + mapped_queues[GfxQueueId].indirect_args_addr = set_base->Address(); + break; + } case PM4ItOpcode::EventWrite: { // const auto* event = reinterpret_cast(header); break; diff --git a/src/video_core/amdgpu/liverpool.h b/src/video_core/amdgpu/liverpool.h index 14284bbc..7f262e1f 100644 --- a/src/video_core/amdgpu/liverpool.h +++ b/src/video_core/amdgpu/liverpool.h @@ -167,7 +167,7 @@ struct Liverpool { static constexpr auto* GetBinaryInfo(const Shader& sh) { const auto* code = sh.template Address(); const auto* bininfo = std::bit_cast(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 ccb_buffer; std::queue submits{}; ComputeProgram cs_state{}; + VAddr indirect_args_addr{}; }; std::array mapped_queues{}; diff --git a/src/video_core/amdgpu/pixel_format.h b/src/video_core/amdgpu/pixel_format.h index 1004ed7d..53d30a7f 100644 --- a/src/video_core/amdgpu/pixel_format.h +++ b/src/video_core/amdgpu/pixel_format.h @@ -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); diff --git a/src/video_core/amdgpu/pm4_cmds.h b/src/video_core/amdgpu/pm4_cmds.h index 5ab233fd..50e4c93a 100644 --- a/src/video_core/amdgpu/pm4_cmds.h +++ b/src/video_core/amdgpu/pm4_cmds.h @@ -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 + T Address() const { + ASSERT(base_index == BaseIndex::DisplayListPatchTable || + base_index == BaseIndex::DrawIndexIndirPatchTable); + return reinterpret_cast(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 diff --git a/src/video_core/amdgpu/resource.h b/src/video_core/amdgpu/resource.h index 8c3b675e..b85a3788 100644 --- a/src/video_core/amdgpu/resource.h +++ b/src/video_core/amdgpu/resource.h @@ -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; } diff --git a/src/video_core/buffer_cache/buffer.cpp b/src/video_core/buffer_cache/buffer.cpp index d112864d..372b6f74 100644 --- a/src/video_core/buffer_cache/buffer.cpp +++ b/src/video_core/buffer_cache/buffer.cpp @@ -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; } + 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().createBufferView({ - .buffer = buffer.buffer, - .format = Vulkan::LiverpoolToVK::SurfaceFormat(dfmt, nfmt), - .offset = offset, - .range = size, - }), + .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(); diff --git a/src/video_core/buffer_cache/buffer.h b/src/video_core/buffer_cache/buffer.h index 87293c76..26d48eae 100644 --- a/src/video_core/buffer_cache/buffer.h +++ b/src/video_core/buffer_cache/buffer.h @@ -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 views; }; diff --git a/src/video_core/buffer_cache/buffer_cache.cpp b/src/video_core/buffer_cache/buffer_cache.cpp index 02d6b2ce..0151f2c1 100644 --- a/src/video_core/buffer_cache/buffer_cache.cpp +++ b/src/video_core/buffer_cache/buffer_cache.cpp @@ -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 BufferCache::ObtainBuffer(VAddr device_addr, u32 size, bool is_written) { +std::pair 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(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) { diff --git a/src/video_core/buffer_cache/buffer_cache.h b/src/video_core/buffer_cache/buffer_cache.h index 2bcc4f0e..a07470b8 100644 --- a/src/video_core/buffer_cache/buffer_cache.h +++ b/src/video_core/buffer_cache/buffer_cache.h @@ -66,7 +66,8 @@ public: u32 BindIndexBuffer(bool& is_indexed, u32 index_offset); /// Obtains a buffer for the specified region. - [[nodiscard]] std::pair ObtainBuffer(VAddr gpu_addr, u32 size, bool is_written); + [[nodiscard]] std::pair 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 ObtainTempBuffer(VAddr gpu_addr, u32 size); diff --git a/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp index 81cf9c02..1d900123 100644 --- a/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp +++ b/src/video_core/renderer_vulkan/vk_compute_pipeline.cpp @@ -13,22 +13,31 @@ 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 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 - : vk::DescriptorType::eUniformBuffer, + .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, }); @@ -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 buffer_views; boost::container::static_vector buffer_infos; boost::container::static_vector image_infos; boost::container::small_vector 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 - : vk::DescriptorType::eUniformBuffer, + .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(image_desc.sgpr_base, image_desc.dword_offset); - 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); + 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({ diff --git a/src/video_core/renderer_vulkan/vk_compute_pipeline.h b/src/video_core/renderer_vulkan/vk_compute_pipeline.h index 5da9dc7e..0132066c 100644 --- a/src/video_core/renderer_vulkan/vk_compute_pipeline.h +++ b/src/video_core/renderer_vulkan/vk_compute_pipeline.h @@ -3,7 +3,7 @@ #pragma once -#include "shader_recompiler/ir/program.h" +#include #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 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 { diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp index 5623e431..1ab65737 100644 --- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp +++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.cpp @@ -19,15 +19,11 @@ namespace Vulkan { GraphicsPipeline::GraphicsPipeline(const Instance& instance_, Scheduler& scheduler_, const GraphicsPipelineKey& key_, vk::PipelineCache pipeline_cache, - std::span programs) + std::span infos, + std::span 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 shader_stages; - shader_stages.emplace_back(vk::PipelineShaderStageCreateInfo{ - .stage = vk::ShaderStageFlagBits::eVertex, - .module = programs[stage]->module, - .pName = "main", - }); + if (infos[stage]) { + shader_stages.emplace_back(vk::PipelineShaderStageCreateInfo{ + .stage = vk::ShaderStageFlagBits::eVertex, + .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,10 +307,20 @@ 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 - : vk::DescriptorType::eUniformBuffer, + .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, }); @@ -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 buffer_infos; + boost::container::static_vector buffer_views; + boost::container::static_vector buffer_infos; boost::container::static_vector image_infos; boost::container::small_vector 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 - : vk::DescriptorType::eUniformBuffer, + .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 tsharps; for (const auto& image_desc : stage->images) { - const auto tsharp = - stage->ReadUd(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) { diff --git a/src/video_core/renderer_vulkan/vk_graphics_pipeline.h b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h index 548e7d45..3e51e652 100644 --- a/src/video_core/renderer_vulkan/vk_graphics_pipeline.h +++ b/src/video_core/renderer_vulkan/vk_graphics_pipeline.h @@ -59,7 +59,8 @@ class GraphicsPipeline { public: explicit GraphicsPipeline(const Instance& instance, Scheduler& scheduler, const GraphicsPipelineKey& key, vk::PipelineCache pipeline_cache, - std::span programs); + std::span stages, + std::span modules); ~GraphicsPipeline(); void BindResources(const Liverpool::Regs& regs, VideoCore::BufferCache& buffer_cache, diff --git a/src/video_core/renderer_vulkan/vk_instance.cpp b/src/video_core/renderer_vulkan/vk_instance.cpp index 19939a32..e1a5cb41 100644 --- a/src/video_core/renderer_vulkan/vk_instance.cpp +++ b/src/video_core/renderer_vulkan/vk_instance.cpp @@ -178,7 +178,7 @@ bool Instance::CreateDevice() { return false; } - boost::container::static_vector enabled_extensions; + boost::container::static_vector 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(); } + if (!maintenance5) { + device_chain.unlink(); + } if (!custom_border_color) { device_chain.unlink(); } diff --git a/src/video_core/renderer_vulkan/vk_instance.h b/src/video_core/renderer_vulkan/vk_instance.h index 4cb4741a..5f985d4a 100644 --- a/src/video_core/renderer_vulkan/vk_instance.h +++ b/src/video_core/renderer_vulkan/vk_instance.h @@ -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; diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp index 139edcf7..7e880657 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp @@ -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 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 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(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(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(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()) { 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 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"); - } - - 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(); - 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(); - 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(instance, scheduler, graphics_key, *pipeline_cache, - programs); -} - -std::unique_ptr 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(); - 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(); - info.pgm_hash = compute_key; - program->pgm = - Shader::TranslateProgram(inst_pool, block_pool, code, std::move(info), profile); - - // Compile IR to SPIR-V - 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(instance, scheduler, *pipeline_cache, compute_key, - it.value().get()); - } catch (const Shader::Exception& e) { - UNREACHABLE_MSG("{}", e.what()); - return nullptr; + const GuestProgram guest_pgm{pgm, stage}; + std::tie(infos[i], modules[i], key.stage_hashes[i]) = + shader_cache->GetProgram(guest_pgm, binding); } } -void PipelineCache::DumpShader(std::span 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); +void PipelineCache::RefreshComputeKey() { + u32 binding{}; + 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 diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.h b/src/video_core/renderer_vulkan/vk_pipeline_cache.h index 8f3b806c..d0eb0c66 100644 --- a/src/video_core/renderer_vulkan/vk_pipeline_cache.h +++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.h @@ -4,9 +4,6 @@ #pragma once #include -#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 code, u64 hash, Shader::Stage stage, std::string_view ext); - - std::unique_ptr CreateGraphicsPipeline(); - std::unique_ptr 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> program_cache; + std::unique_ptr shader_cache; tsl::robin_map> compute_pipelines; tsl::robin_map> graphics_pipelines; - std::array programs{}; - Shader::Profile profile{}; + std::array infos{}; + std::array modules{}; GraphicsPipelineKey graphics_key{}; u64 compute_key{}; - Common::ObjectPool inst_pool; - Common::ObjectPool block_pool; }; } // namespace Vulkan diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index f22f19ea..9231c510 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -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{}; diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.h b/src/video_core/renderer_vulkan/vk_rasterizer.h index a151ebc2..34f6ae72 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.h +++ b/src/video_core/renderer_vulkan/vk_rasterizer.h @@ -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(); diff --git a/src/video_core/renderer_vulkan/vk_shader_cache.cpp b/src/video_core/renderer_vulkan/vk_shader_cache.cpp new file mode 100644 index 00000000..76255712 --- /dev/null +++ b/src/video_core/renderer_vulkan/vk_shader_cache.cpp @@ -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 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 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 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 diff --git a/src/video_core/renderer_vulkan/vk_shader_cache.h b/src/video_core/renderer_vulkan/vk_shader_cache.h new file mode 100644 index 00000000..191e1b08 --- /dev/null +++ b/src/video_core/renderer_vulkan/vk_shader_cache.h @@ -0,0 +1,156 @@ +// SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +#pragma once + +#include +#include +#include +#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 bitset{}; + boost::container::small_vector buffers; + boost::container::small_vector tex_buffers; + boost::container::small_vector 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 modules; + + explicit Program(const Shader::Info& info_) : info{info_} {} +}; + +struct GuestProgram { + Shader::Stage stage; + std::span user_data; + std::span 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 GetProgram(const GuestProgram& pgm, + u32& binding); + +private: + void DumpShader(std::span code, u64 hash, Shader::Stage stage, size_t perm_idx, + std::string_view ext); + vk::ShaderModule CompileModule(Shader::Info& info, std::span 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 program_cache; + Common::ObjectPool inst_pool; + Common::ObjectPool block_pool; + Common::ObjectPool program_pool; +}; + +} // namespace Vulkan diff --git a/src/video_core/texture_cache/image_view.cpp b/src/video_core/texture_cache/image_view.cpp index 00c3833f..bcdc11ad 100644 --- a/src/video_core/texture_cache/image_view.cpp +++ b/src/video_core/texture_cache/image_view.cpp @@ -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); } diff --git a/src/video_core/texture_cache/texture_cache.cpp b/src/video_core/texture_cache/texture_cache.cpp index cae12422..3354a8ec 100644 --- a/src/video_core/texture_cache/texture_cache.cpp +++ b/src/video_core/texture_cache/texture_cache.cpp @@ -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. diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index 8af68424..31b1e393 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -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); diff --git a/src/video_core/texture_cache/tile_manager.cpp b/src/video_core/texture_cache/tile_manager.cpp index 0bed5adc..8b022762 100644 --- a/src/video_core/texture_cache/tile_manager.cpp +++ b/src/video_core/texture_cache/tile_manager.cpp @@ -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: