Compare commits

..
Author SHA1 Message Date
nmzik 6aaa533296 gpu: queue GDS release interrupts behind earlier interrupts 2026-10-23 22:22:39 +02:00
nmzik dc40478e7b gpu: deliver RELEASE_MEM interrupts with the slice flush instead of their own 2026-10-23 22:14:55 +02:00
nmzik 56a82bcd41 launcher: draw the game grid as cards over the selected game's background 2026-10-23 21:24:32 +02:00
nmzik bf5dd62eea gpu: decode RELEASE_MEM directly from the driver's native packet 2026-10-23 21:11:07 +02:00
nmzik 227f12ee4b gpu: remove the PS4-era constant engine, marker NOPs and EVENT_WRITE_EOP/EOS packets 2026-10-23 20:42:41 +02:00
nmzik 246d511568 graphics: count occlusion queries on the host GPU behind --occlusion-queries 2026-10-23 19:48:53 +02:00
nmzik ba0e93daac libs: implement JpegDec and JpegEnc 2026-10-23 19:48:53 +02:00
Senaxx 77a8f30758 shader: read S_MEMREALTIME from the host device clock shifted to about 100 MHz 2026-10-23 19:48:53 +02:00
nmzik 1ffaaa96b1 kernel: share the direct-memory span walk; unchecked release frees the parts around gaps 2026-10-23 19:48:53 +02:00
Senaxx 706aa030ad shader: extract 9-bit NGG primitive vertex indices, leaving the edge flags out 2026-10-23 19:48:53 +02:00
nmzik ecc30a1f2d shader: decode S_FLBIT_I32 and S_FLBIT_I32_I64 with a 4-op signed leading-bit search 2026-10-23 19:48:53 +02:00
nmzik 3688412dca revert stub 2026-10-23 19:48:53 +02:00
nmzik 45f53236e8 hle: implement sceDeviceServiceGetEventState (My First Gran Turismo) 2026-10-23 19:48:53 +02:00
Felipe 9dcff95e80 shader: fix crash from packed ancillary used as DPP old operand (#1769)
shader: emit plain EXEC selects for DPP writes that reach every active lane (Street Fighter 6)
2026-10-23 09:51:32 +02:00
nmzik 719c44310a graphics: mark image heap resources as sampled reads 2026-10-23 08:08:56 +02:00
nmzik 7bb618d41a docs: show Cyberpunk 2077 and MLB The Show 23 in README screenshots 2026-10-23 07:57:57 +02:00
nmzik 87e815b250 net: accept a NULL user agent in sceHttpCreateTemplate (Cyberpunk 2077) 2026-10-23 07:42:28 +02:00
nmzik f9e14084d6 graphics: set the heap flag in ImageHeapResource 2026-10-23 07:22:33 +02:00
nmzik 38c884e8c1 graphics: keep image heap descriptors of unmapped textures null 2026-10-23 07:17:57 +02:00
nmzik 27dc7c281f graphics: keep bounded image heaps resident in a bindless descriptor set 2026-10-23 07:17:57 +02:00
88 changed files with 3963 additions and 1385 deletions

No files matched your search

+21
View File
@@ -541,6 +541,19 @@ add_executable(png_enc_tests EXCLUDE_FROM_ALL "${KYTY_TESTS_DIR}/PngEncTests.cpp
target_link_libraries(png_enc_tests common fmt::fmt)
target_include_directories(png_enc_tests PRIVATE ${inc_headers})
add_executable(jpeg_dec_tests EXCLUDE_FROM_ALL "${KYTY_TESTS_DIR}/JpegDecTests.cpp")
target_link_libraries(jpeg_dec_tests common fmt::fmt)
target_include_directories(jpeg_dec_tests PRIVATE ${inc_headers})
add_executable(jpeg_enc_tests EXCLUDE_FROM_ALL "${KYTY_TESTS_DIR}/JpegEncTests.cpp")
target_link_libraries(jpeg_enc_tests common fmt::fmt)
target_include_directories(jpeg_enc_tests PRIVATE ${inc_headers})
add_executable(occlusion_queries_tests EXCLUDE_FROM_ALL
"${KYTY_TESTS_DIR}/OcclusionQueriesTests.cpp")
target_link_libraries(occlusion_queries_tests common fmt::fmt Vulkan::Headers)
target_include_directories(occlusion_queries_tests PRIVATE ${inc_headers})
add_executable(ces_tests EXCLUDE_FROM_ALL "${KYTY_TESTS_DIR}/CesTests.cpp")
target_link_libraries(ces_tests common fmt::fmt SDL3::SDL3-static)
target_include_directories(ces_tests PRIVATE ${inc_headers})
@@ -652,6 +665,9 @@ if(BUILD_TESTING)
add_test(NAME host_input COMMAND $<TARGET_FILE:host_input_tests>)
add_test(NAME save_data_memory COMMAND $<TARGET_FILE:save_data_memory_tests>)
add_test(NAME png_enc COMMAND $<TARGET_FILE:png_enc_tests>)
add_test(NAME jpeg_dec COMMAND $<TARGET_FILE:jpeg_dec_tests>)
add_test(NAME jpeg_enc COMMAND $<TARGET_FILE:jpeg_enc_tests>)
add_test(NAME occlusion_queries COMMAND $<TARGET_FILE:occlusion_queries_tests>)
add_test(NAME ces COMMAND $<TARGET_FILE:ces_tests>)
add_test(NAME save_data_dialog COMMAND $<TARGET_FILE:save_data_dialog_tests>)
add_test(NAME http_uri_parse COMMAND $<TARGET_FILE:http_uri_parse_tests>)
@@ -678,6 +694,8 @@ if(BUILD_TESTING)
COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --gpu-command-lane-only)
add_test(NAME wait_until_safe_for_rendering
COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --wait-until-safe-for-rendering-only)
add_test(NAME pm4_flip
COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --flip-only)
add_test(NAME depth_bits
COMMAND $<TARGET_FILE:shader_recompiler_compute_tests> --depth-bits-only)
add_test(NAME msaa_storage
@@ -742,6 +760,9 @@ if(BUILD_TESTING)
host_input_tests
save_data_memory_tests
png_enc_tests
jpeg_dec_tests
jpeg_enc_tests
occlusion_queries_tests
ces_tests
save_data_dialog_tests
http_uri_parse_tests
+10
View File
@@ -46,6 +46,16 @@ or graphical glitches, so please include the version you tested when reporting a
## Screenshots
<table align="center">
<tr>
<td align="center">
<strong>Cyberpunk 2077</strong><br>
<img src="docs/screenshots/ps5-07.png" width="300" alt="Cyberpunk 2077 running in KytyPS5">
</td>
<td align="center">
<strong>MLB The Show 23</strong><br>
<img src="docs/screenshots/ps5-08.png" width="300" alt="MLB The Show 23 running in KytyPS5">
</td>
</tr>
<tr>
<td align="center">
<strong>Astro Bot</strong><br>
Binary file not shown.

After

Width:  |  Height:  |  Size: 1.5 MiB

Binary file not shown.

After

Width:  |  Height:  |  Size: 2.0 MiB

+4
View File
@@ -197,6 +197,10 @@ bool TessellationEnabled() {
return g_config->tessellation_enabled;
}
bool OcclusionQueriesEnabled() {
return g_config->occlusion_queries_enabled;
}
bool TrophyEnabled() {
return g_config->trophy_enabled;
}
+2
View File
@@ -88,6 +88,7 @@ struct ConfigOptions {
bool readback_tiled_images = false;
bool sync_raw_image_buffers = false;
bool tessellation_enabled = false;
bool occlusion_queries_enabled = false;
bool trophy_enabled = true;
uint32_t trophy_duration_seconds = 6;
bool perf_overlay_enabled = false;
@@ -154,6 +155,7 @@ bool ReadbackLinearImagesEnabled();
bool ReadbackTiledImagesEnabled();
bool SyncRawImageBuffersEnabled();
bool TessellationEnabled();
bool OcclusionQueriesEnabled();
bool TrophyEnabled();
uint32_t GetTrophyDurationSeconds();
bool PlayGoHackEnabled();
@@ -50,13 +50,6 @@ private:
class CommandProcessor {
public:
struct FlipInfo {
int handle = 0;
int index = 0;
int flip_mode = 0;
int64_t flip_arg = 0;
};
CommandProcessor(RenderContext& renderer, int interrupt_event_id)
: m_renderer(renderer), m_interrupt_event_id(interrupt_event_id) {}
~CommandProcessor() = default;
@@ -90,40 +83,17 @@ public:
void DrawIndirectMulti(uint32_t data_offset, uint32_t max_count_or_count,
const volatile uint32_t* count_addr, uint32_t stride_in_bytes,
IndirectDrawRegisters registers, uint32_t draw_initiator, bool indexed);
void WriteAtEndOfPipe32(uint32_t cache_policy, uint32_t event_write_dest,
uint32_t eop_event_type, uint32_t cache_action, uint32_t event_index,
uint32_t event_write_source, void* dst_gpu_addr, uint32_t value,
uint32_t interrupt_selector, uint32_t interrupt_context_id = 0);
void WriteAtEndOfPipe64(uint32_t cache_policy, uint32_t event_write_dest,
uint32_t eop_event_type, uint32_t cache_action, uint32_t event_index,
uint32_t event_write_source, void* dst_gpu_addr, uint64_t value,
uint32_t interrupt_selector, uint32_t interrupt_context_id = 0);
void Flip();
void Flip(void* dst_gpu_addr, uint32_t value);
void FlipWithInterrupt(uint32_t eop_event_type, uint32_t cache_action, void* dst_gpu_addr,
uint32_t value);
void ReleaseMem(uint32_t data_sel, uint32_t interrupt_selector, void* address, uint64_t data,
uint32_t interrupt_context_id);
void Flip(int handle, int index, int flip_mode, int64_t flip_arg);
void PrepareCpuFlip(uint64_t request_id);
void SynchronizeGpu();
void TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id);
void DispatchDirect(uint32_t thread_group_x, uint32_t thread_group_y, uint32_t thread_group_z,
uint32_t mode);
void DispatchIndirect(uint64_t args_addr, uint32_t mode);
void WaitFlipDone(uint32_t video_out_handle, uint32_t display_buffer_index);
void TriggerEvent(uint32_t event_type, uint32_t event_index, uint64_t event_address = 0);
void SetUserDataMarker(HW::UserSgprType type) { m_user_data_marker = type; }
[[nodiscard]] HW::UserSgprType GetUserDataMarker() const { return m_user_data_marker; }
void ResetDeCe();
void SetCeComplete(bool complete) { m_ce_complete = complete; }
void WaitCe();
void WaitDeDiff(uint32_t diff);
void WaitForRewind(bool valid);
void IncrementDe();
void IncrementCe();
void WriteConstRam(uint32_t offset, const uint32_t* src, uint32_t dw_num);
void DumpConstRam(uint32_t* dst, uint32_t offset, uint32_t dw_num);
template <typename T>
void WaitRegMem(uint32_t func, const T* addr, T ref, T mask, uint32_t poll, uint32_t wait_op);
@@ -140,18 +110,11 @@ public:
Pm4ProcessResult Process(Pm4Execution& execution, std::span<const uint32_t> commands);
void ProcessIndirectBuffer(std::span<const uint32_t> commands, bool chain);
void SetFlip(const FlipInfo& flip) { m_flip = flip; }
[[nodiscard]] uint64_t GetSubmitId() const { return m_submit_id; }
void SetSubmitId(uint64_t submit_id) { m_submit_id = submit_id; }
[[nodiscard]] bool IsAsyncComputeQueue() const { return m_interrupt_event_id >= 0x20; }
private:
template <typename T>
void WriteAtEndOfPipe(uint32_t cache_policy, uint32_t event_write_dest, uint32_t eop_event_type,
uint32_t cache_action, uint32_t event_index, uint32_t event_write_source,
void* dst_gpu_addr, T value, uint32_t interrupt_selector,
uint32_t interrupt_context_id);
void ProcessPm4(Pm4Execution& execution);
void SuspendPm4();
uint32_t NumInstances();
@@ -159,30 +122,22 @@ private:
CommandScheduler& GetScheduler() const { return m_renderer.GetCommandScheduler(); }
CommandBuffer& CurrentBuffer() { return GetScheduler().Current(); }
RenderContext& m_renderer;
HW::Context m_ctx;
HW::Context m_saved_ctx;
bool m_context_state_pushed = false;
HW::UserConfig m_ucfg;
HW::Shader m_sh_ctx;
HW::UserSgprType m_user_data_marker = HW::UserSgprType::Unknown;
uint32_t m_index_type_and_size = 0;
uint32_t m_index_buffer_size = 0;
uint64_t m_index_base_addr = 0;
uint64_t m_draw_indirect_args_base_addr = 0;
uint64_t m_dispatch_indirect_args_base_addr = 0;
RenderContext& m_renderer;
HW::Context m_ctx;
HW::Context m_saved_ctx;
bool m_context_state_pushed = false;
HW::UserConfig m_ucfg;
HW::Shader m_sh_ctx;
uint32_t m_index_type_and_size = 0;
uint32_t m_index_buffer_size = 0;
uint64_t m_index_base_addr = 0;
uint64_t m_draw_indirect_args_base_addr = 0;
uint64_t m_dispatch_indirect_args_base_addr = 0;
// Persistent draw state. Indirect draws leave the address of their last instance count,
// read only when a later draw inherits it.
uint32_t m_num_instances = 1;
uint64_t m_num_instances_addr = 0;
uint32_t m_de_count = 0;
uint32_t m_ce_count = 0;
bool m_ce_complete = false;
uint32_t m_const_ram[0x3000] = {0};
FlipInfo m_flip;
const int m_interrupt_event_id;
uint64_t m_submit_id = 0;
uint64_t m_synthetic_occlusion_counter = 0;
@@ -199,10 +199,9 @@ constexpr auto MakeOpcodeDispatchTable() {
g_cp_op_func[Pm4::IT_INDIRECT_BUFFER] = CpOpIndirectBuffer;
g_cp_op_func[Pm4::IT_COPY_DATA] = CpOpCopyData;
g_cp_op_func[Pm4::IT_EVENT_WRITE] = CpOpEventWrite;
g_cp_op_func[Pm4::IT_EVENT_WRITE_EOP] = CpOpEventWriteEop;
g_cp_op_func[Pm4::IT_EVENT_WRITE_EOS] = CpOpEventWriteEos;
g_cp_op_func[Pm4::IT_DMA_DATA] = CpOpDmaData;
g_cp_op_func[Pm4::IT_ACQUIRE_MEM] = CpOpAcquireMem;
g_cp_op_func[Pm4::IT_RELEASE_MEM] = CpOpReleaseMem;
g_cp_op_func[Pm4::IT_REWIND] = CpOpRewind;
g_cp_op_func[Pm4::IT_SET_CONTEXT_REG] = CpOpSetContextReg;
g_cp_op_func[Pm4::IT_SET_SH_REG] = CpOpSetShaderReg;
@@ -211,12 +210,6 @@ constexpr auto MakeOpcodeDispatchTable() {
g_cp_op_func[Pm4::IT_SET_CONTEXT_REG_INDIRECT] = CpOpIndirectCxRegs;
g_cp_op_func[Pm4::IT_SET_SH_REG_INDIRECT] = CpOpIndirectShRegs;
g_cp_op_func[Pm4::IT_SET_UCONFIG_REG_INDIRECT] = CpOpIndirectUcRegs;
g_cp_op_func[Pm4::IT_WRITE_CONST_RAM] = CpOpWriteConstRam;
g_cp_op_func[Pm4::IT_DUMP_CONST_RAM] = CpOpDumpConstRam;
g_cp_op_func[Pm4::IT_INCREMENT_CE_COUNTER] = CpOpIncrementCeCounter;
g_cp_op_func[Pm4::IT_INCREMENT_DE_COUNTER] = CpOpIncrementDeCounter;
g_cp_op_func[Pm4::IT_WAIT_ON_CE_COUNTER] = CpOpWaitOnCeCounter;
g_cp_op_func[Pm4::IT_WAIT_ON_DE_COUNTER_DIFF] = CpOpWaitOnDeCounterDiff;
g_cp_op_func[Pm4::IT_GET_LOD_STATS] = CpOpGetLodStats;
g_cp_op_func[Pm4::IT_WAIT_REG_MEM_64] = CpOpWaitRegMem64;
@@ -231,7 +224,6 @@ constexpr auto MakeCustomOpcodeDispatchTable() {
g_cp_op_custom_func[Pm4::R_PUSH_MARKER] = CpOpPushMarker;
g_cp_op_custom_func[Pm4::R_POP_MARKER] = CpOpPopMarker;
g_cp_op_custom_func[Pm4::R_FLIP] = CpOpFlip;
g_cp_op_custom_func[Pm4::R_RELEASE_MEM] = CpOpReleaseMem;
g_cp_op_custom_func[Pm4::R_CONTEXT_STATE] = CpOpContextState;
return g_cp_op_custom_func;
@@ -133,22 +133,15 @@ uint32_t CpOpWriteData(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, u
uint32_t CpOpIndirectBuffer(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpCopyData(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpEventWrite(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpEventWriteEop(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpEventWriteEos(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpDmaData(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpAcquireMem(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpReleaseMem(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpSetContextReg(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpSetShaderReg(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpSetUconfigReg(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpIndirectCxRegs(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpIndirectShRegs(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpIndirectUcRegs(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpWriteConstRam(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpDumpConstRam(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpIncrementCeCounter(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpIncrementDeCounter(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpWaitOnCeCounter(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpWaitOnDeCounterDiff(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpGetLodStats(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpDispatchReset(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpWaitRegMem32(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
@@ -157,7 +150,6 @@ uint32_t CpOpPushMarker(CommandProcessor&, uint32_t, const uint32_t*, uint32_t,
uint32_t CpOpPopMarker(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpWaitRegMem64(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpFlip(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
uint32_t CpOpReleaseMem(CommandProcessor&, uint32_t, const uint32_t*, uint32_t, uint32_t);
} // namespace Libs::Graphics
@@ -54,34 +54,12 @@ namespace Libs::Graphics {
namespace {
constexpr uint32_t GcrGl2MetadataInvalidate = 1u << 1u;
constexpr uint32_t GcrGl0VectorInvalidate = 1u << 2u;
constexpr uint32_t GcrGl1Invalidate = 1u << 3u;
constexpr uint32_t GcrGl2Unshared = 1u << 4u;
constexpr uint32_t GcrGl2Invalidate = 1u << 8u;
constexpr uint32_t GcrGl2Writeback = 1u << 9u;
constexpr uint32_t GcrOrder012 = 1u << 10u;
constexpr uint32_t GcrOrder210 = 2u << 10u;
constexpr uint32_t GcrKnownMask = GcrGl2MetadataInvalidate | GcrGl0VectorInvalidate |
GcrGl1Invalidate | GcrGl2Unshared | GcrGl2Invalidate |
GcrGl2Writeback | GcrOrder012 | GcrOrder210;
constexpr uint32_t RegisterSelectorMask = 0x70000000u;
constexpr uint32_t RegisterSelectorMask = 0x70000000u;
constexpr uint32_t NormalizeRegisterOffset(uint32_t raw_offset) {
return raw_offset & ~RegisterSelectorMask;
}
uint32_t ReleaseMemCacheActionFromGcr(uint32_t gcr_cntl) {
return ((gcr_cntl & GcrGl2Writeback) != 0 ? 0x38u : 0x00u);
}
void LogUnknownReleaseMemGcr(uint32_t gcr_cntl) {
const uint32_t unknown = (gcr_cntl & ~GcrKnownMask);
if (unknown != 0) {
LOGF("\t warning: release_mem uses unknown GCR bits: 0x%04" PRIx32 "\n", unknown);
}
}
} // namespace
static HW::BlendControl DecodeBlendControl(uint32_t value) {
@@ -1068,9 +1046,8 @@ KYTY_HW_SH_PARSER(HwShSetCsUserSgpr) {
auto write_num = UserSgprWriteNum(slot, reg_num, 16u, "cs");
for (uint32_t i = 0; i < write_num; i++) {
cp.GetShCtx().SetCsUserSgpr(slot + i, buffer[i], cp.GetUserDataMarker());
cp.GetShCtx().SetCsUserSgpr(slot + i, buffer[i]);
}
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
return reg_num;
}
@@ -1085,9 +1062,8 @@ KYTY_HW_SH_PARSER(HwShSetPsUserSgpr) {
auto write_num = UserSgprWriteNum(slot, reg_num, 32u, "ps");
for (uint32_t i = 0; i < write_num; i++) {
cp.GetShCtx().SetPsUserSgpr(slot + i, buffer[i], cp.GetUserDataMarker());
cp.GetShCtx().SetPsUserSgpr(slot + i, buffer[i]);
}
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
return reg_num;
}
@@ -1102,9 +1078,8 @@ KYTY_HW_SH_PARSER(HwShSetGsUserSgpr) {
auto write_num = UserSgprWriteNum(slot, reg_num, 32u, "gs");
for (uint32_t i = 0; i < write_num; i++) {
cp.GetShCtx().SetGsUserSgpr(slot + i, buffer[i], cp.GetUserDataMarker());
cp.GetShCtx().SetGsUserSgpr(slot + i, buffer[i]);
}
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
return reg_num;
}
@@ -1119,9 +1094,8 @@ KYTY_HW_SH_PARSER(HwShSetHsUserSgpr) {
auto write_num = UserSgprWriteNum(slot, reg_num, 32u, "hs");
for (uint32_t i = 0; i < write_num; i++) {
cp.GetShCtx().SetHsUserSgpr(slot + i, buffer[i], cp.GetUserDataMarker());
cp.GetShCtx().SetHsUserSgpr(slot + i, buffer[i]);
}
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
return reg_num;
}
@@ -1700,24 +1674,6 @@ KYTY_CP_OP_PARSER(CpOpContextState) {
return packet_size_dw - 1u;
}
KYTY_CP_OP_PARSER(CpOpDumpConstRam) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0038300);
auto offset = buffer[0];
auto dw_num = buffer[1];
auto* dst = reinterpret_cast<uint32_t*>(buffer[2] | (static_cast<uint64_t>(buffer[3]) << 32u));
EXIT_NOT_IMPLEMENTED(dw_num >= 0x3000);
EXIT_NOT_IMPLEMENTED(offset > 0xbffc);
EXIT_NOT_IMPLEMENTED((offset & 0x3u) != 0);
cp.DumpConstRam(dst, offset, dw_num);
return 4;
}
KYTY_CP_OP_PARSER(CpOpEventWrite) {
KYTY_PROFILER_FUNCTION();
@@ -1738,90 +1694,16 @@ KYTY_CP_OP_PARSER(CpOpEventWrite) {
return packet_size_dw - 1u;
}
KYTY_CP_OP_PARSER(CpOpEventWriteEop) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0044700);
uint32_t cache_policy = (buffer[0] >> 25u) & 0x3u;
uint32_t event_write_dest = ((buffer[0] >> 23u) & 0x10u) | ((buffer[2] >> 16u) & 0x01u);
uint32_t eop_event_type = (buffer[0]) & 0x3fu;
uint32_t cache_action = (buffer[0] >> 12u) & 0x3fu;
uint32_t event_index = (buffer[0] >> 8u) & 0x7u;
uint32_t event_write_source = ((buffer[2] >> 29u) & 0x7u);
uint32_t interrupt_selector = (buffer[2] >> 24u) & 0x7u;
auto* dst_gpu_addr =
reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2] & 0xffffu) << 32u));
uint64_t value = (buffer[3] | (static_cast<uint64_t>(buffer[4]) << 32u));
cp.WriteAtEndOfPipe64(cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
event_write_source, dst_gpu_addr, value, interrupt_selector);
return 5;
}
KYTY_CP_OP_PARSER(CpOpEventWriteEos) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0034802);
uint32_t cache_policy = (buffer[0] >> 25u) & 0x3u;
uint32_t event_write_dest = 0;
uint32_t eop_event_type = (buffer[0]) & 0x3fu;
uint32_t cache_action = (buffer[0] >> 12u) & 0x3fu;
uint32_t event_index = (buffer[0] >> 8u) & 0x7u;
uint32_t event_write_source = ((buffer[2] >> 29u) & 0x7u);
uint32_t interrupt_selector = (buffer[2] >> 24u) & 0x7u;
auto* dst_gpu_addr =
reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2] & 0xffffu) << 32u));
uint32_t value = buffer[3];
cp.WriteAtEndOfPipe32(cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
event_write_source, dst_gpu_addr, value, interrupt_selector);
return 4;
}
KYTY_CP_OP_PARSER(CpOpFlip) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xc004105c);
CommandProcessor::FlipInfo f;
f.handle = static_cast<int>(buffer[0]);
f.index = static_cast<int>(buffer[1]);
f.flip_mode = static_cast<int>(buffer[2]);
f.flip_arg = static_cast<int64_t>(buffer[3] | (static_cast<uint64_t>(buffer[4]) << 32u));
cp.SetFlip(f);
cp.Flip();
cp.Flip(static_cast<int>(buffer[0]), static_cast<int>(buffer[1]), static_cast<int>(buffer[2]),
static_cast<int64_t>(buffer[3] | (static_cast<uint64_t>(buffer[4]) << 32u)));
return 5;
}
KYTY_CP_OP_PARSER(CpOpIncrementCeCounter) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0008400);
EXIT_NOT_IMPLEMENTED(buffer[0] != 1);
cp.IncrementCe();
return 1;
}
KYTY_CP_OP_PARSER(CpOpIncrementDeCounter) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0008500);
EXIT_NOT_IMPLEMENTED(buffer[0] != 0);
cp.IncrementDe();
return 1;
}
KYTY_CP_OP_PARSER(CpOpIndexType) {
KYTY_PROFILER_FUNCTION();
@@ -2103,50 +1985,12 @@ KYTY_CP_OP_PARSER(CpOpIndirectUcRegs) {
return KYTY_PM4_LEN(cmd_id) - 1u;
}
KYTY_CP_OP_PARSER(CpOpMarker) {
KYTY_PROFILER_FUNCTION();
uint32_t len_dw = ((cmd_id >> 16u) & 0x3fffu);
switch (buffer[0]) {
case 0x68750004: cp.SetUserDataMarker(HW::UserSgprType::Vsharp); break;
case 0x6875000d: cp.SetUserDataMarker(HW::UserSgprType::Region); break;
case 0x68750777: {
cp.Flip();
break;
}
case 0x68750778: {
auto* addr =
reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2]) << 32u));
uint32_t value = buffer[3];
cp.Flip(addr, value);
break;
}
case 0x68750781: {
auto* addr =
reinterpret_cast<void*>(buffer[1] | (static_cast<uint64_t>(buffer[2]) << 32u));
uint32_t value = buffer[3];
uint32_t eop_event_type = buffer[4];
uint32_t cache_action = buffer[5];
cp.FlipWithInterrupt(eop_event_type, cache_action, addr, value);
break;
}
// NOP payloads are caller-owned data; only complete known marker tags execute.
default: break;
}
return len_dw + 1;
}
KYTY_CP_OP_PARSER(CpOpNop) {
KYTY_PROFILER_FUNCTION();
auto r = KYTY_PM4_R(cmd_id);
if (r == Pm4::R_ZERO) {
if ((buffer[0] & 0xffff0000u) == 0x68750000u) {
return CpOpMarker(cp, cmd_id, buffer, dw, num_dw);
}
return KYTY_PM4_LEN(cmd_id) - 1;
}
@@ -2202,108 +2046,14 @@ KYTY_CP_OP_PARSER(CpOpPushMarker) {
KYTY_CP_OP_PARSER(CpOpReleaseMem) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xc0061060);
EXIT_NOT_IMPLEMENTED(cmd_id != 0xc0064900);
uint32_t cache_policy = (buffer[0] >> 25u) & 0x3u;
uint32_t eop_event_type = buffer[0] & 0x3fu;
uint32_t event_index = (buffer[0] >> 8u) & 0x7u;
uint32_t gcr_cntl = (buffer[0] >> 12u) & 0xfffu;
uint32_t release_dst = (buffer[1] >> 16u) & 0x3u;
uint32_t event_write_dest = 0;
uint32_t data_sel = (buffer[1] >> 29u) & 0x7u;
uint32_t interrupt_selector = (buffer[1] >> 24u) & 0x7u;
auto* dst_gpu_addr =
reinterpret_cast<void*>(buffer[2] | (static_cast<uint64_t>(buffer[3]) << 32u));
uint64_t value = buffer[4] | (static_cast<uint64_t>(buffer[5]) << 32u);
uint32_t interrupt_context_id = buffer[6] & 0x07ffffffu;
constexpr uint32_t ReleaseMemDstMemory = 0;
constexpr uint32_t ReleaseMemDstTcL2 = 1;
EXIT_NOT_IMPLEMENTED(release_dst != ReleaseMemDstMemory && release_dst != ReleaseMemDstTcL2);
EXIT_NOT_IMPLEMENTED(data_sel != 0 && data_sel != 1 && data_sel != 2 && data_sel != 3 &&
data_sel != 5);
LogUnknownReleaseMemGcr(gcr_cntl);
const bool gl2_writeback = ((gcr_cntl & GcrGl2Writeback) != 0);
auto trigger_interrupt = [&]() {
switch (interrupt_selector) {
case 0x00:
case 0x03: break;
case 0x01:
case 0x02:
case 0x04:
cp.TriggerEopEventAtEndOfPipe(interrupt_context_id);
cp.BufferFlush();
break;
default: EXIT("unknown release_mem interrupt selector\n");
}
};
// Cache actions only order GPU work, which the renderer's resource barriers already order.
if (data_sel == 0 || interrupt_selector == 4 ||
(release_dst == ReleaseMemDstMemory && dst_gpu_addr == nullptr)) {
trigger_interrupt();
return 7;
}
auto cache_action = ReleaseMemCacheActionFromGcr(gcr_cntl);
cache_policy = 0;
if (data_sel == 1) {
eop_event_type = 0x2f;
event_index = 6;
auto event_source = 2u;
cp.WriteAtEndOfPipe32(cache_policy, event_write_dest, eop_event_type, cache_action,
event_index, event_source, dst_gpu_addr, static_cast<uint32_t>(value),
interrupt_selector, interrupt_context_id);
if (interrupt_selector == 0x01 || interrupt_selector == 0x02) {
cp.BufferFlush();
}
return 7;
}
if (data_sel == 5) {
if (gl2_writeback) {
cache_action = 0;
}
eop_event_type = 0x2f;
event_index = 6;
auto event_source = 1u;
cp.WriteAtEndOfPipe32(cache_policy, event_write_dest, eop_event_type, cache_action,
event_index, event_source, dst_gpu_addr, static_cast<uint32_t>(value),
interrupt_selector, interrupt_context_id);
if (interrupt_selector == 0x01) {
cp.BufferFlush();
}
return 7;
}
const bool keep_native_event_index =
(eop_event_type == 0x04 && cache_action == 0x00) ||
(eop_event_type == 0x2f && event_index == 0x06 && cache_action == 0x38);
if (data_sel == 2) {
if (!keep_native_event_index) {
event_index = 0;
}
} else {
EXIT_IF(data_sel != 3);
if (!keep_native_event_index) {
event_index = 0;
}
data_sel = 4;
}
cp.WriteAtEndOfPipe64(cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
data_sel, dst_gpu_addr, value, interrupt_selector, interrupt_context_id);
// The event type, cache actions and destination cache only order GPU work, which the renderer's
// resource barriers already order.
cp.ReleaseMem(
(buffer[1] >> 29u) & 0x7u, (buffer[1] >> 24u) & 0x7u,
reinterpret_cast<void*>((buffer[2] & ~0x3u) | (static_cast<uint64_t>(buffer[3]) << 32u)),
buffer[4] | (static_cast<uint64_t>(buffer[5]) << 32u), buffer[6] & 0x07ffffffu);
return 7;
}
@@ -2524,42 +2274,6 @@ KYTY_CP_OP_PARSER(CpOpWaitRegMem64) {
return CpOpWaitRegMemSized<uint64_t>(cp, cmd_id, buffer);
}
KYTY_CP_OP_PARSER(CpOpWaitOnCeCounter) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0008600);
EXIT_NOT_IMPLEMENTED(buffer[0] != 1);
cp.WaitCe();
return 1;
}
KYTY_CP_OP_PARSER(CpOpWaitOnDeCounterDiff) {
KYTY_PROFILER_FUNCTION();
EXIT_NOT_IMPLEMENTED(cmd_id != 0xC0008800);
cp.WaitDeDiff(buffer[0]);
return 1;
}
KYTY_CP_OP_PARSER(CpOpWriteConstRam) {
KYTY_PROFILER_FUNCTION();
auto dw_num = (cmd_id >> 16u) & 0x3fffu;
auto offset = buffer[0];
EXIT_NOT_IMPLEMENTED(dw_num >= 0x3000);
EXIT_NOT_IMPLEMENTED(offset > 0xbffc);
EXIT_NOT_IMPLEMENTED((offset & 0x3u) != 0);
cp.WriteConstRam(offset, buffer + 1, dw_num);
return 1 + dw_num;
}
KYTY_CP_OP_PARSER(CpOpWriteData) {
KYTY_PROFILER_FUNCTION();
@@ -3519,9 +3233,7 @@ void GraphicsInitJmpTablesShIndirect() {
for (uint32_t slot = 0; slot < 16; slot++) {
g_hw_sh_indirect_func[Pm4::COMPUTE_USER_DATA_0 + slot] = [](KYTY_HW_SH_INDIRECT_ARGS) {
auto sgpr = cmd_offset - Pm4::COMPUTE_USER_DATA_0;
cp.GetShCtx().SetCsUserSgpr(sgpr, value, cp.GetUserDataMarker());
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
cp.GetShCtx().SetCsUserSgpr(cmd_offset - Pm4::COMPUTE_USER_DATA_0, value);
};
}
for (uint32_t slot = 0; slot < 4; slot++) {
@@ -3532,21 +3244,15 @@ void GraphicsInitJmpTablesShIndirect() {
for (uint32_t slot = 0; slot < 32; slot++) {
g_hw_sh_indirect_func[Pm4::SPI_SHADER_USER_DATA_PS_0 + slot] =
[](KYTY_HW_SH_INDIRECT_ARGS) {
auto sgpr = cmd_offset - Pm4::SPI_SHADER_USER_DATA_PS_0;
cp.GetShCtx().SetPsUserSgpr(sgpr, value, cp.GetUserDataMarker());
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
cp.GetShCtx().SetPsUserSgpr(cmd_offset - Pm4::SPI_SHADER_USER_DATA_PS_0, value);
};
g_hw_sh_indirect_func[Pm4::SPI_SHADER_USER_DATA_GS_0 + slot] =
[](KYTY_HW_SH_INDIRECT_ARGS) {
auto sgpr = cmd_offset - Pm4::SPI_SHADER_USER_DATA_GS_0;
cp.GetShCtx().SetGsUserSgpr(sgpr, value, cp.GetUserDataMarker());
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
cp.GetShCtx().SetGsUserSgpr(cmd_offset - Pm4::SPI_SHADER_USER_DATA_GS_0, value);
};
g_hw_sh_indirect_func[Pm4::SPI_SHADER_USER_DATA_HS_0 + slot] =
[](KYTY_HW_SH_INDIRECT_ARGS) {
auto sgpr = cmd_offset - Pm4::SPI_SHADER_USER_DATA_HS_0;
cp.GetShCtx().SetHsUserSgpr(sgpr, value, cp.GetUserDataMarker());
cp.SetUserDataMarker(HW::UserSgprType::Unknown);
cp.GetShCtx().SetHsUserSgpr(cmd_offset - Pm4::SPI_SHADER_USER_DATA_HS_0, value);
};
}
for (uint32_t slot = 0; slot < 4; slot++) {
+83 -359
View File
@@ -127,16 +127,14 @@ void GuestGpu::SendCommandSync(Common::UniqueFunction<void>&& command) {
done.acquire();
}
void GuestGpu::Submit(std::span<const uint32_t> draw_commands,
std::span<const uint32_t> constant_commands) {
if (draw_commands.empty()) {
void GuestGpu::Submit(std::span<const uint32_t> commands) {
if (commands.empty()) {
return;
}
Submission submission;
submission.type = SubmissionType::Graphics;
submission.queue_id = 0;
submission.commands = draw_commands;
submission.constant_commands = constant_commands;
submission.type = SubmissionType::Graphics;
submission.queue_id = 0;
submission.commands = commands;
Enqueue(std::move(submission));
}
@@ -200,11 +198,8 @@ void CommandProcessor::Reset() {
m_num_instances = 1;
m_num_instances_addr = 0;
m_predicate_skip = false;
m_user_data_marker = HW::UserSgprType::Unknown;
m_draw_indirect_args_base_addr = 0;
m_dispatch_indirect_args_base_addr = 0;
std::memset(m_const_ram, 0, sizeof(m_const_ram));
}
void CommandProcessor::ApplyContextStateOperation(ContextStateOperation operation) {
@@ -237,6 +232,8 @@ void CommandProcessor::BufferInit() {
void CommandProcessor::BufferFlush() {
GetScheduler().Flush();
// The occlusion counters reach a guest that polls them without waiting for a later fence.
m_renderer.GetOcclusionQueries().Publish();
}
void CommandProcessor::BufferWait() {
@@ -244,47 +241,12 @@ void CommandProcessor::BufferWait() {
GetScheduler().Finish();
}
void CommandProcessor::ResetDeCe() {
m_de_count = 0;
m_ce_count = 0;
m_ce_complete = false;
}
void CommandProcessor::WaitCe() {
if (m_ce_count <= m_de_count && !m_ce_complete) {
SuspendPm4();
}
}
void CommandProcessor::WaitDeDiff(uint32_t diff) {
EXIT_IF(m_de_count > m_ce_count);
if (m_ce_count - m_de_count >= diff) {
SuspendPm4();
}
}
void CommandProcessor::WaitForRewind(bool valid) {
if (!valid) {
SuspendPm4();
}
}
void CommandProcessor::IncrementDe() {
m_de_count++;
}
void CommandProcessor::IncrementCe() {
m_ce_count++;
}
void CommandProcessor::WriteConstRam(uint32_t offset, const uint32_t* src, uint32_t dw_num) {
memcpy(m_const_ram + offset / 4, src, static_cast<size_t>(dw_num) * 4);
}
void CommandProcessor::DumpConstRam(uint32_t* dst, uint32_t offset, uint32_t dw_num) {
memcpy(dst, m_const_ram + offset / 4, static_cast<size_t>(dw_num) * 4);
}
bool TestWaitRegMemValue(uint64_t value, uint64_t ref, uint64_t mask, uint32_t func) {
switch (func) {
case 0: return true;
@@ -539,64 +501,34 @@ bool GuestGpu::Process(Submission& submission) {
if (first_slice) {
submission.started = true;
cp.SetSubmitId(++m_submit_id);
cp.ResetDeCe();
cp.SetFlip({});
}
cp.BufferInit();
bool complete = true;
switch (submission.type) {
case SubmissionType::Graphics: {
bool progressed = false;
submission.constant_complete |= submission.constant_commands.empty();
for (;;) {
bool round_progress = false;
if (!submission.constant_complete) {
submission.constant_complete =
cp.Process(submission.constant_execution, submission.constant_commands) ==
Pm4ProcessResult::Complete;
round_progress |= submission.constant_execution.MadeProgress();
}
cp.SetCeComplete(submission.constant_complete);
if (!submission.command_complete) {
submission.command_complete =
cp.Process(submission.command_execution, submission.commands) ==
Pm4ProcessResult::Complete;
round_progress |= submission.command_execution.MadeProgress();
}
progressed |= round_progress;
complete = submission.command_complete && submission.constant_complete;
if (complete || !round_progress) {
break;
}
}
if (progressed) {
cp.BufferFlush();
}
break;
}
case SubmissionType::Compute: {
const auto num_dw = static_cast<uint32_t>(submission.commands.size());
const auto* buffer = submission.commands.data();
static uint32_t compute_batch_log_count = 0;
if (first_slice && num_dw <= 128 && compute_batch_log_count++ < 32) {
LOGF("compute direct batch: data=0x%016" PRIx64 ", num_dw=%" PRIu32 "\n",
reinterpret_cast<uint64_t>(buffer), num_dw);
for (uint32_t i = 0; i < std::min<uint32_t>(num_dw, 16); i++) {
LOGF("\t compute[%02" PRIu32 "] = 0x%08" PRIx32 "\n", i, buffer[i]);
}
}
case SubmissionType::Compute:
if (first_slice) {
const auto num_dw = static_cast<uint32_t>(submission.commands.size());
const auto* buffer = submission.commands.data();
static uint32_t compute_batch_log_count = 0;
if (num_dw <= 128 && compute_batch_log_count++ < 32) {
LOGF("compute direct batch: data=0x%016" PRIx64 ", num_dw=%" PRIu32 "\n",
reinterpret_cast<uint64_t>(buffer), num_dw);
for (uint32_t i = 0; i < std::min<uint32_t>(num_dw, 16); i++) {
LOGF("\t compute[%02" PRIu32 "] = 0x%08" PRIx32 "\n", i, buffer[i]);
}
}
GraphicsDbgDumpDcb("cc", num_dw, buffer);
}
complete = cp.Process(submission.command_execution, submission.commands) ==
Pm4ProcessResult::Complete;
if (submission.command_execution.MadeProgress()) {
[[fallthrough]];
case SubmissionType::Graphics:
complete =
cp.Process(submission.execution, submission.commands) == Pm4ProcessResult::Complete;
if (submission.execution.MadeProgress()) {
cp.BufferFlush();
}
break;
}
case SubmissionType::FlipPreparation:
cp.PrepareCpuFlip(submission.flip_request_id);
break;
@@ -699,18 +631,6 @@ void CommandProcessor::ProcessPm4(Pm4Execution& execution) {
opcode, KYTY_PM4_R(packet_header), packet_dw,
reinterpret_cast<uint64_t>(packet), packet_header);
}
if (opcode == Pm4::IT_NOP && KYTY_PM4_R(packet_header) == Pm4::R_RELEASE_MEM &&
packet_dw >= 7) {
static std::atomic<uint32_t> log_count {0};
if (log_count.fetch_add(1) < 128) {
const auto dst = packet[3] | (static_cast<uint64_t>(packet[4]) << 32u);
const auto val = packet[5] | (static_cast<uint64_t>(packet[6]) << 32u);
LOGF("\t predicated skip: R_RELEASE_MEM dst=0x%016" PRIx64
", value=0x%016" PRIx64 ", action=0x%08" PRIx32
", gcr/data/int=0x%08" PRIx32 "\n",
dst, val, packet[1], packet[2]);
}
}
cursor.offset_dw += packet_dw;
execution.m_made_progress = true;
continue;
@@ -803,6 +723,7 @@ void CommandProcessor::SetPredication(uint32_t condition, uint32_t op, uint32_t
return;
case 0x01: {
EXIT_NOT_IMPLEMENTED(address == nullptr);
m_renderer.GetOcclusionQueries().Publish();
// One begin/end pair per DB; bit 63 marks each counter ready.
constexpr uint64_t ready_bit = 1ull << 63u;
const auto* results = reinterpret_cast<const volatile uint64_t*>(address);
@@ -1008,216 +929,68 @@ void CommandProcessor::DrawIndexAuto(DrawAutoArgs args) {
}
void CommandProcessor::WaitFlipDone(uint32_t video_out_handle, uint32_t display_buffer_index) {
// The SDK stalls only this command processor, not the host command lane or other queues.
// Stall only this command processor, not the host command lane or other queues.
if (!m_renderer.GetVideoOut().IsSafeForRendering(static_cast<int>(video_out_handle),
static_cast<int>(display_buffer_index))) {
SuspendPm4();
}
}
template <typename T>
void CommandProcessor::WriteAtEndOfPipe(uint32_t cache_policy, uint32_t event_write_dest,
uint32_t eop_event_type, uint32_t cache_action,
uint32_t event_index, uint32_t event_write_source,
void* dst_gpu_addr, T value, uint32_t interrupt_selector,
uint32_t interrupt_context_id) {
static_assert(sizeof(T) == sizeof(uint32_t) || sizeof(T) == sizeof(uint64_t));
auto& command = CurrentBuffer();
void CommandProcessor::ReleaseMem(uint32_t data_sel, uint32_t interrupt_selector, void* address,
uint64_t data, uint32_t interrupt_context_id) {
if (GraphicsRunDebugDumpEnabled()) {
const auto bits = static_cast<unsigned>(sizeof(T) * 8u);
const auto log_width = static_cast<int>(sizeof(T) * 2u);
LOGF("CommandProcessor::WriteAtEndOfPipe%u()\n"
"\t cache_policy = 0x%08" PRIx32 "\n"
"\t event_write_dest = 0x%08" PRIx32 "\n"
"\t eop_event_type = 0x%08" PRIx32 "\n"
"\t cache_action = 0x%08" PRIx32 "\n"
"\t event_index = 0x%08" PRIx32 "\n"
"\t event_write_source = 0x%08" PRIx32 "\n"
"\t interrupt_selector = 0x%08" PRIx32 "\n"
"\t interrupt_context = 0x%08" PRIx32 "\n"
"\t dst_gpu_addr = 0x%016" PRIx64 "\n"
"\t value = 0x%0*" PRIx64 "\n",
bits, cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
event_write_source, interrupt_selector, interrupt_context_id,
reinterpret_cast<uint64_t>(dst_gpu_addr), log_width, static_cast<uint64_t>(value));
LOGF("CommandProcessor::ReleaseMem()\n"
"\t data_sel = %" PRIu32 "\n"
"\t interrupt_selector = %" PRIu32 "\n"
"\t address = 0x%016" PRIx64 "\n"
"\t data = 0x%016" PRIx64 "\n"
"\t interrupt_context = 0x%08" PRIx32 "\n",
data_sel, interrupt_selector, reinterpret_cast<uint64_t>(address), data,
interrupt_context_id);
}
EXIT_NOT_IMPLEMENTED(cache_policy != 0x00000000);
EXIT_NOT_IMPLEMENTED(event_write_dest != 0x00000000);
// QueuedInterrupt: kNone, kOnly, kAfterWriteConfirm, kDataAfterWriteConfirm, kSendIntCtx.
EXIT_NOT_IMPLEMENTED(interrupt_selector > 4);
const bool interrupt = (interrupt_selector != 0 && interrupt_selector != 3);
// kSendIntCtx has no label, and kOnly writes none on graphics queues.
if (address == nullptr || interrupt_selector == 4 ||
(interrupt_selector == 1 && !IsAsyncComputeQueue())) {
data_sel = 0;
}
if (data_sel == 0 && !interrupt) {
return;
}
bool with_interrupt = false;
switch (interrupt_selector) {
case 0x00:
case 0x03: with_interrupt = false; break;
case 0x01:
if (!IsAsyncComputeQueue()) {
TriggerEopEventAtEndOfPipe(interrupt_context_id);
return;
}
with_interrupt = true;
// Dumped occlusion counters are written before any end-of-pipe signal reaches the guest.
m_renderer.GetOcclusionQueries().Publish();
// EndOfPipeData: kNone, kImmediate32b, kImmediate64b, kRefClk; 5 reads GDS.
switch (data_sel) {
case 0: break;
case 1: {
const auto value = static_cast<uint32_t>(data);
std::memcpy(address, &value, sizeof(value));
break;
case 0x02: with_interrupt = true; break;
default: EXIT("unknown interrupt_selector\n");
}
auto write32 = [&] {
auto* dst = static_cast<uint32_t*>(dst_gpu_addr);
auto data = static_cast<uint32_t>(value);
EXIT_IF(dst == nullptr);
(void)command.Handle();
std::memcpy(dst, &data, sizeof(data));
if (with_interrupt) {
TriggerEopEventAtEndOfPipe(interrupt_context_id);
}
};
switch (event_write_source) {
case 0x01:
if constexpr (sizeof(T) == sizeof(uint32_t)) {
if (eop_event_type == 0x2f && cache_action == 0x00 && event_index == 0x06) {
auto* dst = static_cast<uint32_t*>(dst_gpu_addr);
SynchronizeGpu();
Sync::ReadGds(*m_renderer.GetBufferCache().GetGdsBuffer(), dst, value & 0xffffu,
value >> 16u);
if (with_interrupt) {
m_renderer.TriggerInterrupt(m_interrupt_event_id, interrupt_context_id);
}
return;
}
} else if (eop_event_type == 0x04 && cache_action == 0x00 && event_index == 0x05) {
write32();
return;
}
case 3: data = Sync::ReadReferenceClock(); [[fallthrough]];
case 2: std::memcpy(address, &data, sizeof(data)); break;
case 5: {
GetScheduler().Finish();
const auto gds = static_cast<uint32_t>(data);
Sync::ReadGds(*m_renderer.GetBufferCache().GetGdsBuffer(),
static_cast<uint32_t*>(address), gds & 0xffffu, gds >> 16u);
break;
case 0x02:
case 0x04:
if constexpr (sizeof(T) == sizeof(uint32_t)) {
if (event_write_source == 0x02 && eop_event_type == 0x2f && event_index == 0x06) {
switch (cache_action) {
case 0x00:
case 0x38: write32(); return;
default: break;
}
}
} else {
if (event_write_source == 0x04) {
value = Sync::ReadReferenceClock();
}
auto write64 = [&] {
auto* dst = static_cast<uint64_t*>(dst_gpu_addr);
EXIT_IF(dst == nullptr);
(void)command.Handle();
std::memcpy(dst, &value, sizeof(value));
if (with_interrupt) {
TriggerEopEventAtEndOfPipe(interrupt_context_id);
}
};
switch (cache_action) {
case 0x00:
switch (eop_event_type) {
case 0x04:
if (event_index == 0x05) {
write64();
return;
}
break;
case 0x14:
case 0x28:
case 0x2f:
if (event_index == 0x00) {
write64();
return;
}
break;
case 0x2b:
case 0x2d:
case 0x30:
if (event_index == 0x00 && !with_interrupt) {
write64();
return;
}
break;
default: break;
}
break;
case 0x38:
switch (eop_event_type) {
case 0x04:
case 0x14:
case 0x28:
if (((eop_event_type == 0x04 || eop_event_type == 0x28) &&
event_index == 0x05) ||
(event_index == 0x00)) {
write64();
return;
}
break;
case 0x2b:
case 0x2d:
if (event_index == 0x00 && !with_interrupt) {
write64();
return;
}
break;
case 0x2f:
if (event_index == 0x06 && !with_interrupt) {
write64();
return;
}
break;
default: break;
}
break;
case 0x3b:
if (eop_event_type == 0x04 && event_index == 0x05 && with_interrupt) {
write64();
return;
}
break;
default: break;
}
}
break;
default: break;
}
default: EXIT("unknown release_mem data selector: %" PRIu32 "\n", data_sel);
}
EXIT("unknown event type\n");
}
void CommandProcessor::WriteAtEndOfPipe32(uint32_t cache_policy, uint32_t event_write_dest,
uint32_t eop_event_type, uint32_t cache_action,
uint32_t event_index, uint32_t event_write_source,
void* dst_gpu_addr, uint32_t value,
uint32_t interrupt_selector,
uint32_t interrupt_context_id) {
WriteAtEndOfPipe(cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
event_write_source, dst_gpu_addr, value, interrupt_selector,
interrupt_context_id);
}
void CommandProcessor::WriteAtEndOfPipe64(uint32_t cache_policy, uint32_t event_write_dest,
uint32_t eop_event_type, uint32_t cache_action,
uint32_t event_index, uint32_t event_write_source,
void* dst_gpu_addr, uint64_t value,
uint32_t interrupt_selector,
uint32_t interrupt_context_id) {
WriteAtEndOfPipe(cache_policy, event_write_dest, eop_event_type, cache_action, event_index,
event_write_source, dst_gpu_addr, value, interrupt_selector,
interrupt_context_id);
}
void CommandProcessor::TriggerEopEventAtEndOfPipe(uint32_t interrupt_context_id) {
(void)CurrentBuffer().Handle();
GetScheduler().DeferPriorityOperation(
[&renderer = m_renderer, event_id = m_interrupt_event_id, interrupt_context_id] {
renderer.TriggerInterrupt(event_id, interrupt_context_id);
});
// GuestGpu::Process's end-of-slice flush submits the work this interrupt waits on.
if (interrupt) {
GetScheduler().DeferPriorityOperation(
[&renderer = m_renderer, event_id = m_interrupt_event_id, interrupt_context_id] {
renderer.TriggerInterrupt(event_id, interrupt_context_id);
});
}
}
void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
@@ -1270,15 +1043,19 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
"\n",
event_index, event_address);
}
if (auto& queries = m_renderer.GetOcclusionQueries(); queries.Enabled()) {
queries.Dump(event_address);
break;
}
static std::once_flag warning_once;
std::call_once(warning_once, [] {
std::printf("Warning: game uses occlusion queries, which are currently treated as "
"always visible; GPU usage may be higher and FPS may be lower.\n");
std::printf("Warning: game uses occlusion queries, which are treated as always "
"visible without --occlusion-queries; GPU usage may be higher and FPS "
"may be lower.\n");
});
// Until host occlusion queries are implemented, publish an always-visible result. The
// PS5 layout contains one interleaved begin/end pair per DB, and bit 63 marks a result
// ready.
// Without host occlusion queries, publish an always-visible result. The PS5 layout
// contains one interleaved begin/end pair per DB, and bit 63 marks a result ready.
constexpr uint64_t ready_bit = 1ull << 63u;
constexpr uint64_t counter_mask = ready_bit - 1u;
auto* results = reinterpret_cast<volatile uint64_t*>(event_address);
@@ -1294,69 +1071,20 @@ void CommandProcessor::TriggerEvent(uint32_t event_type, uint32_t event_index,
}
}
void CommandProcessor::Flip() {
void CommandProcessor::Flip(int handle, int index, int flip_mode, int64_t flip_arg) {
if (GraphicsRunDebugDumpEnabled()) {
LOGF("CommandProcessor::Flip()\n");
}
m_renderer.GetOcclusionQueries().Publish();
auto& command = CurrentBuffer();
auto request = Sync::PrepareVideoOutFlip(command, m_flip.handle, m_flip.index, m_flip.flip_mode,
m_flip.flip_arg);
auto request = Sync::PrepareVideoOutFlip(command, handle, index, flip_mode, flip_arg);
(void)command.Handle();
GetScheduler().DeferPriorityOperation(
[&renderer = m_renderer, request] { renderer.GetVideoOut().CompleteFlip(request); });
GetScheduler().Flush();
}
void CommandProcessor::Flip(void* dst_gpu_addr, uint32_t value) {
auto& command = CurrentBuffer();
if (GraphicsRunDebugDumpEnabled()) {
LOGF("CommandProcessor::Flip()\n"
"\t dst_gpu_addr = 0x%016" PRIx64 "\n"
"\t value = 0x%08" PRIx32 "\n",
reinterpret_cast<uint64_t>(dst_gpu_addr), value);
}
EXIT_IF(dst_gpu_addr == nullptr);
std::memcpy(dst_gpu_addr, &value, sizeof(value));
auto request = Sync::PrepareVideoOutFlip(command, m_flip.handle, m_flip.index, m_flip.flip_mode,
m_flip.flip_arg);
(void)command.Handle();
GetScheduler().DeferPriorityOperation(
[&renderer = m_renderer, request] { renderer.GetVideoOut().CompleteFlip(request); });
GetScheduler().Flush();
}
void CommandProcessor::FlipWithInterrupt(uint32_t eop_event_type, uint32_t cache_action,
void* dst_gpu_addr, uint32_t value) {
auto& command = CurrentBuffer();
if (GraphicsRunDebugDumpEnabled()) {
LOGF("CommandProcessor::FlipWithInterrupt()\n"
"\t eop_event_type = 0x%08" PRIx32 "\n"
"\t cache_action = 0x%08" PRIx32 "\n"
"\t dst_gpu_addr = 0x%016" PRIx64 "\n"
"\t value = 0x%08" PRIx32 "\n",
eop_event_type, cache_action, reinterpret_cast<uint64_t>(dst_gpu_addr), value);
}
if (eop_event_type != 0x00000004 || cache_action != 0x00000038) {
EXIT("unknown event type\n");
}
EXIT_IF(dst_gpu_addr == nullptr);
std::memcpy(dst_gpu_addr, &value, sizeof(value));
auto request = Sync::PrepareVideoOutFlip(command, m_flip.handle, m_flip.index, m_flip.flip_mode,
m_flip.flip_arg);
(void)command.Handle();
GetScheduler().DeferPriorityOperation(
[&renderer = m_renderer, event_id = m_interrupt_event_id, request] {
renderer.GetVideoOut().CompleteFlip(request);
renderer.TriggerInterrupt(event_id, 0);
});
GetScheduler().Flush();
}
void CommandProcessor::PrepareCpuFlip(uint64_t request_id) {
auto& command = CurrentBuffer();
if (g_current_processor != nullptr) {
@@ -1374,10 +1102,6 @@ void CommandProcessor::PrepareCpuFlip(uint64_t request_id) {
GetScheduler().Flush();
}
void CommandProcessor::SynchronizeGpu() {
GetScheduler().Finish();
}
bool GuestGpu::IsGpuThread() noexcept {
return g_gpu_thread;
}
+5 -10
View File
@@ -35,8 +35,7 @@ public:
void Wake();
// Submitted command memory is borrowed and must remain valid until GPU execution completes.
void Submit(std::span<const uint32_t> draw_commands,
std::span<const uint32_t> constant_commands);
void Submit(std::span<const uint32_t> commands);
void SubmitCompute(uint32_t queue, std::span<const uint32_t> commands);
void SubmitFlipPreparation(uint64_t request_id);
// Insert an ordered graphics drain; only a previous suspend point can block the caller.
@@ -60,14 +59,10 @@ private:
SubmissionType type = SubmissionType::Graphics;
uint32_t queue_id = 0;
std::span<const uint32_t> commands;
std::span<const uint32_t> constant_commands;
Pm4Execution command_execution;
Pm4Execution constant_execution;
bool started = false;
bool command_complete = false;
bool constant_complete = false;
bool blocked = false;
uint64_t flip_request_id = 0;
Pm4Execution execution;
bool started = false;
bool blocked = false;
uint64_t flip_request_id = 0;
};
void Enqueue(Submission submission);
+5 -29
View File
@@ -597,14 +597,10 @@ struct ShaderRegisters {
}
};
enum class UserSgprType { Unknown, Region, Vsharp };
struct UserSgprInfo {
static constexpr int SGPRS_MAX = 32;
uint32_t value[SGPRS_MAX] = {0};
UserSgprType type[SGPRS_MAX] = {};
uint32_t count = 0;
uint32_t value[SGPRS_MAX] = {0};
};
struct VertexShaderInfo {
@@ -1056,30 +1052,10 @@ public:
void SetCsShader(const CsStageRegisters& cs_regs) { m_cs.cs_regs = cs_regs; }
void SetCsWaveSize(uint8_t wave_size) { m_cs.cs_regs.wave_size = wave_size; }
void SetPsUserSgpr(uint32_t id, uint32_t value, UserSgprType type) {
m_ps.ps_user_sgpr.value[id] = value;
m_ps.ps_user_sgpr.type[id] = type;
m_ps.ps_user_sgpr.count =
((id + 1) > m_ps.ps_user_sgpr.count ? (id + 1) : m_ps.ps_user_sgpr.count);
}
void SetCsUserSgpr(uint32_t id, uint32_t value, UserSgprType type) {
m_cs.cs_user_sgpr.value[id] = value;
m_cs.cs_user_sgpr.type[id] = type;
m_cs.cs_user_sgpr.count =
((id + 1) > m_cs.cs_user_sgpr.count ? (id + 1) : m_cs.cs_user_sgpr.count);
}
void SetHsUserSgpr(uint32_t id, uint32_t value, UserSgprType type) {
m_vs.hs_user_sgpr.value[id] = value;
m_vs.hs_user_sgpr.type[id] = type;
m_vs.hs_user_sgpr.count =
((id + 1) > m_vs.hs_user_sgpr.count ? (id + 1) : m_vs.hs_user_sgpr.count);
}
void SetGsUserSgpr(uint32_t id, uint32_t value, UserSgprType type) {
m_vs.gs_user_sgpr.value[id] = value;
m_vs.gs_user_sgpr.type[id] = type;
m_vs.gs_user_sgpr.count =
((id + 1) > m_vs.gs_user_sgpr.count ? (id + 1) : m_vs.gs_user_sgpr.count);
}
void SetPsUserSgpr(uint32_t id, uint32_t value) { m_ps.ps_user_sgpr.value[id] = value; }
void SetCsUserSgpr(uint32_t id, uint32_t value) { m_cs.cs_user_sgpr.value[id] = value; }
void SetHsUserSgpr(uint32_t id, uint32_t value) { m_vs.hs_user_sgpr.value[id] = value; }
void SetGsUserSgpr(uint32_t id, uint32_t value) { m_vs.gs_user_sgpr.value[id] = value; }
[[nodiscard]] const PixelShaderInfo& GetPs() const { return m_ps; }
[[nodiscard]] const VertexShaderInfo& GetVs() const { return m_vs; }
-10
View File
@@ -25,7 +25,6 @@ constexpr auto MakeRegisterNames() {
names[R_POP_MARKER] = "R_POP_MARKER";
names[R_WRITE_DATA] = "R_WRITE_DATA";
names[R_FLIP] = "R_FLIP";
names[R_RELEASE_MEM] = "R_RELEASE_MEM";
names[R_DMA_DATA] = "R_DMA_DATA";
names[R_CONTEXT_STATE] = "R_CONTEXT_STATE";
return names;
@@ -54,7 +53,6 @@ constexpr auto MakeOpcodeNames() {
names[IT_DRAW_INDIRECT_MULTI] = "IT_DRAW_INDIRECT_MULTI";
names[IT_DRAW_INDEX_AUTO] = "IT_DRAW_INDEX_AUTO";
names[IT_NUM_INSTANCES] = "IT_NUM_INSTANCES";
names[IT_INDIRECT_BUFFER_CNST] = "IT_INDIRECT_BUFFER_CNST";
names[IT_DRAW_INDEX_OFFSET_2] = "IT_DRAW_INDEX_OFFSET_2";
names[IT_WRITE_DATA] = "IT_WRITE_DATA";
names[IT_MEM_SEMAPHORE] = "IT_MEM_SEMAPHORE";
@@ -66,8 +64,6 @@ constexpr auto MakeOpcodeNames() {
names[IT_PFP_SYNC_ME] = "IT_PFP_SYNC_ME";
names[IT_SURFACE_SYNC] = "IT_SURFACE_SYNC";
names[IT_EVENT_WRITE] = "IT_EVENT_WRITE";
names[IT_EVENT_WRITE_EOP] = "IT_EVENT_WRITE_EOP";
names[IT_EVENT_WRITE_EOS] = "IT_EVENT_WRITE_EOS";
names[IT_RELEASE_MEM] = "IT_RELEASE_MEM";
names[IT_DMA_DATA] = "IT_DMA_DATA";
names[IT_ACQUIRE_MEM] = "IT_ACQUIRE_MEM";
@@ -80,12 +76,6 @@ constexpr auto MakeOpcodeNames() {
names[IT_SET_QUEUE_REG] = "IT_SET_QUEUE_REG";
names[IT_SET_UCONFIG_REG] = "IT_SET_UCONFIG_REG";
names[IT_SET_UCONFIG_REG_INDEX] = "IT_SET_UCONFIG_REG_INDEX";
names[IT_WRITE_CONST_RAM] = "IT_WRITE_CONST_RAM";
names[IT_DUMP_CONST_RAM] = "IT_DUMP_CONST_RAM";
names[IT_INCREMENT_CE_COUNTER] = "IT_INCREMENT_CE_COUNTER";
names[IT_INCREMENT_DE_COUNTER] = "IT_INCREMENT_DE_COUNTER";
names[IT_WAIT_ON_CE_COUNTER] = "IT_WAIT_ON_CE_COUNTER";
names[IT_WAIT_ON_DE_COUNTER_DIFF] = "IT_WAIT_ON_DE_COUNTER_DIFF";
names[IT_DISPATCH_DRAW_PREAMBLE] = "IT_DISPATCH_DRAW_PREAMBLE";
names[IT_DISPATCH_DRAW] = "IT_DISPATCH_DRAW";
names[IT_GET_LOD_STATS] = "IT_GET_LOD_STATS";
-10
View File
@@ -40,7 +40,6 @@ constexpr uint32_t IT_INDEX_TYPE = 0x2A;
constexpr uint32_t IT_DRAW_INDIRECT_MULTI = 0x2C;
constexpr uint32_t IT_DRAW_INDEX_AUTO = 0x2D;
constexpr uint32_t IT_NUM_INSTANCES = 0x2F;
constexpr uint32_t IT_INDIRECT_BUFFER_CNST = 0x33;
constexpr uint32_t IT_DRAW_INDEX_OFFSET_2 = 0x35;
constexpr uint32_t IT_WRITE_DATA = 0x37;
constexpr uint32_t IT_MEM_SEMAPHORE = 0x39;
@@ -53,8 +52,6 @@ constexpr uint32_t IT_CP_DMA = 0x41;
constexpr uint32_t IT_PFP_SYNC_ME = 0x42;
constexpr uint32_t IT_SURFACE_SYNC = 0x43;
constexpr uint32_t IT_EVENT_WRITE = 0x46;
constexpr uint32_t IT_EVENT_WRITE_EOP = 0x47;
constexpr uint32_t IT_EVENT_WRITE_EOS = 0x48;
constexpr uint32_t IT_RELEASE_MEM = 0x49;
constexpr uint32_t IT_DMA_DATA = 0x50;
constexpr uint32_t IT_ACQUIRE_MEM = 0x58;
@@ -67,12 +64,6 @@ constexpr uint32_t IT_SET_SH_REG = 0x76;
constexpr uint32_t IT_SET_QUEUE_REG = 0x78;
constexpr uint32_t IT_SET_UCONFIG_REG = 0x79;
constexpr uint32_t IT_SET_UCONFIG_REG_INDEX = 0x7A;
constexpr uint32_t IT_WRITE_CONST_RAM = 0x81;
constexpr uint32_t IT_DUMP_CONST_RAM = 0x83;
constexpr uint32_t IT_INCREMENT_CE_COUNTER = 0x84;
constexpr uint32_t IT_INCREMENT_DE_COUNTER = 0x85;
constexpr uint32_t IT_WAIT_ON_CE_COUNTER = 0x86;
constexpr uint32_t IT_WAIT_ON_DE_COUNTER_DIFF = 0x88;
constexpr uint32_t IT_DISPATCH_DRAW = 0x8D;
constexpr uint32_t IT_GET_LOD_STATS = 0x8E;
constexpr uint32_t IT_WAIT_REG_MEM_64 = 0x93;
@@ -88,7 +79,6 @@ constexpr uint32_t R_PUSH_MARKER = 0x0B;
constexpr uint32_t R_POP_MARKER = 0x0C;
constexpr uint32_t R_WRITE_DATA = 0x15;
constexpr uint32_t R_FLIP = 0x17;
constexpr uint32_t R_RELEASE_MEM = 0x18;
constexpr uint32_t R_DMA_DATA = 0x19;
constexpr uint32_t R_CONTEXT_STATE = 0x1A;
+4
View File
@@ -36,6 +36,8 @@ struct GraphicContext {
bool provoking_vertex_last_enabled = false;
bool maintenance8_enabled = false;
bool conditional_rendering_enabled = false;
bool occlusion_queries_enabled = false;
bool occlusion_query_precise_enabled = false;
bool supports_block_texel_view = false;
bool mesh_shader_enabled = false;
vk::PhysicalDeviceMeshShaderPropertiesEXT mesh_shader_properties = {};
@@ -43,6 +45,8 @@ struct GraphicContext {
uint32_t min_subgroup_size = 0;
uint32_t max_subgroup_size = 0;
uint32_t max_push_descriptors = 0;
uint32_t device_clock_period_ps = 0;
vk::DescriptorSetLayout image_heap_layout = nullptr;
vk::ShaderStageFlags required_subgroup_size_stages = {};
Common::Mutex queue_mutex;
uint32_t queue_family = static_cast<uint32_t>(-1);
+10 -1
View File
@@ -229,6 +229,7 @@ void TextureCache::UnregisterImage(ImageId id) {
if (!image.registered) {
return;
}
image.HeapChanged();
UntrackImage(id);
ImagePageTable::PageRange pages {};
if (!ImagePageTable::TryGetPageRange(image.info.data.address, image.info.data.size, pages)) {
@@ -378,6 +379,7 @@ void TextureCache::UntrackImageHead(ImageId id) {
const auto address = Common::AlignDown(begin + TRACKER_PAGE_SIZE, TRACKER_PAGE_SIZE);
const auto size = address - begin;
image.track_addr = address;
image.HeapChanged();
if (image.track_addr == image.track_addr_end) {
MarkAsMaybeDirty(id, image);
}
@@ -395,6 +397,7 @@ void TextureCache::UntrackImageTail(ImageId id) {
const auto address = Common::AlignDown(end, TRACKER_PAGE_SIZE);
const auto size = end - address;
image.track_addr_end = address;
image.HeapChanged();
if (image.track_addr == image.track_addr_end) {
MarkAsMaybeDirty(id, image);
}
@@ -795,7 +798,8 @@ TextureCache::OverlapResult TextureCache::ResolveOverlap(const ImageInfo& reques
auto& cached = *owner;
const auto current_tick = m_scheduler.CurrentTick();
const bool safe_to_delete =
current_tick - std::min(current_tick, cached.tick_accessed_last) > NumFramesBeforeRemoval;
current_tick - std::min(current_tick, cached.tick_accessed_last) > NumFramesBeforeRemoval &&
cached.heap_refs.load(std::memory_order_relaxed) == 0u;
const uint32_t requested_block = requested.bytes_per_block * requested.samples;
const uint32_t cached_block = cached.info.bytes_per_block * cached.info.samples;
@@ -2015,6 +2019,11 @@ void TextureCache::RunGarbageCollector() {
if (owner == nullptr || !owner->registered || owner->depth_id) {
continue;
}
if (owner->heap_refs.load(std::memory_order_relaxed) != 0u) {
// Resident heaps sample it without touching it.
TouchImage(*owner);
continue;
}
if (owner->IsGpuModified()) {
const bool safe = owner->SafeToDownload();
if (safe && owner->info.IsTiled()) {
+1
View File
@@ -191,6 +191,7 @@ private:
friend struct TextureCacheTestAccess;
friend class BufferCache;
friend class RenderExecutor;
friend class ImageHeap;
friend bool BcAlias::Matches(const ImageInfo&, const ImageInfo&);
friend bool BcAlias::Publish(TextureCache&, Image&);
friend void BcAlias::Resolve(TextureCache&, ImageId, ImageId);
@@ -212,6 +212,13 @@ Image::Barriers Image::GetBarriers(vk::ImageLayout destinat
barriers.push_back(barrier);
}
// Resident heaps sample their images in the read state; only leaving it invalidates them.
if (std::ranges::any_of(barriers, [&](const vk::ImageMemoryBarrier2& barrier) {
return barrier.oldLayout == ReadLayout() &&
barrier.srcAccessMask == vk::AccessFlagBits2::eShaderRead;
})) {
HeapChanged();
}
state = {destination_stage, destination_access, destination_layout};
return barriers;
}
+22 -1
View File
@@ -7,6 +7,7 @@
#include "graphics/host_gpu/graphicContext.h"
#include "graphics/host_gpu/renderer/image/imageInfo.h"
#include <atomic>
#include <compare>
#include <limits>
#include <optional>
@@ -71,6 +72,7 @@ public:
void MarkCpuDirty() {
m_cpu_dirty = true;
m_maybe_cpu_dirty = false;
HeapChanged();
}
[[nodiscard]] bool IsCpuDirty() const { return m_cpu_dirty || m_maybe_cpu_dirty; }
@@ -80,6 +82,7 @@ public:
if (!IsCpuDirty()) {
m_maybe_cpu_hash = HashGuestEdges();
m_maybe_cpu_dirty = true;
HeapChanged();
}
}
[[nodiscard]] bool ResolveMaybeCpuHash(uint64_t hash) {
@@ -104,7 +107,10 @@ public:
void ClearGpuModified() noexcept { m_gpu_modified = false; }
[[nodiscard]] bool IsBufferModified() const noexcept { return m_buffer_modified; }
void MarkBufferModified() noexcept { m_buffer_modified = true; }
void MarkBufferModified() noexcept {
m_buffer_modified = true;
HeapChanged();
}
void ClearBufferModified() noexcept { m_buffer_modified = false; }
[[nodiscard]] bool Overlaps(uint64_t address, uint64_t size,
@@ -120,6 +126,19 @@ public:
return backing.image == nullptr ? 0 : Common::AlignUp(info.data.size, 1024);
}
[[nodiscard]] uint64_t HashGuestEdges() const;
[[nodiscard]] vk::ImageLayout ReadLayout() const noexcept {
return info.IsDepth() ? vk::ImageLayout::eDepthStencilReadOnlyOptimal
: vk::ImageLayout::eShaderReadOnlyOptimal;
}
// Resident image heaps sample an image without rediscovering it until one changes its
// contents, tracking, identity or layout.
void HeapChanged() const noexcept {
if (heap_refs.load(std::memory_order_relaxed) != 0u) {
heap_epoch.fetch_add(1u, std::memory_order_relaxed);
}
}
static inline std::atomic<uint64_t> heap_epoch {0};
ImageInfo info;
VulkanImage backing;
@@ -137,6 +156,8 @@ public:
ImageSubresourceRange stencil_subresources;
uint64_t tick_accessed_last = 0;
size_t lru_id = 0;
// Resident image heap elements that select this image.
std::atomic<uint32_t> heap_refs {0};
private:
friend struct ImageTestAccess;
@@ -0,0 +1,82 @@
#include "graphics/host_gpu/renderer/occlusionQueries.h"
#include "graphics/host_gpu/graphicContext.h"
#include "graphics/host_gpu/renderer/commandScheduler.h"
#include <algorithm>
namespace Libs::Graphics {
// Queries no dump has read yet; past them a draw counts one sample.
constexpr uint32_t QUERY_COUNT = 16384;
OcclusionQueries::OcclusionQueries(GraphicContext& graphics, CommandScheduler& scheduler)
: m_scheduler(scheduler) {
if (!graphics.occlusion_queries_enabled) {
return;
}
m_device = graphics.device;
const vk::QueryPoolCreateInfo info {.queryType = vk::QueryType::eOcclusion,
.queryCount = QUERY_COUNT};
RequireVulkanSuccess(m_device.createQueryPool(&info, nullptr, &m_pool), "vkCreateQueryPool");
m_device.resetQueryPool(m_pool, 0, QUERY_COUNT);
m_results.resize(QUERY_COUNT);
if (graphics.occlusion_query_precise_enabled) {
m_control = vk::QueryControlFlagBits::ePrecise;
}
}
OcclusionQueries::~OcclusionQueries() {
if (m_pool) {
m_device.destroyQueryPool(m_pool, nullptr);
}
}
int32_t OcclusionQueries::BeginQuery(vk::CommandBuffer command) {
if (m_begun - m_read == QUERY_COUNT) {
m_uncounted++;
return -1;
}
const auto query = static_cast<uint32_t>(m_begun++ % QUERY_COUNT);
command.beginQuery(m_pool, query, m_control);
return static_cast<int32_t>(query);
}
void OcclusionQueries::Dump(uint64_t address) {
// Begins dump to the results, ends 8 bytes into them.
if ((address & 8u) == 0) {
m_open++;
} else if (m_open != 0) {
m_open--;
}
m_dumps.push_back({address, m_scheduler.CurrentTick(), m_begun, m_uncounted});
}
void OcclusionQueries::PublishDumps() {
if (const auto last = m_dumps.back().queries; last != m_read) {
m_scheduler.Wait(m_dumps.back().tick);
// The unread queries wrap around the pool at most once.
for (uint64_t read = m_read; read < last;) {
const auto first = static_cast<uint32_t>(read % QUERY_COUNT);
const auto count =
static_cast<uint32_t>(std::min<uint64_t>(last - read, QUERY_COUNT - first));
RequireVulkanSuccess(
m_device.getQueryPoolResults(m_pool, first, count, count * sizeof(uint64_t),
m_results.data() + first, sizeof(uint64_t),
vk::QueryResultFlagBits::e64),
"vkGetQueryPoolResults");
m_device.resetQueryPool(m_pool, first, count);
read += count;
}
}
for (const auto& dump: m_dumps) {
for (; m_read < dump.queries; m_read++) {
m_samples += m_results[m_read % QUERY_COUNT];
}
WriteZPassCounters(reinterpret_cast<volatile uint64_t*>(dump.address),
m_samples + dump.uncounted, (dump.address & 8u) == 0);
}
m_dumps.clear();
}
} // namespace Libs::Graphics
@@ -0,0 +1,85 @@
#ifndef EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_OCCLUSIONQUERIES_H_
#define EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_OCCLUSIONQUERIES_H_
#include "common/common.h"
#include "graphics/host_gpu/vulkanCommon.h"
#include <cstdint>
#include <vector>
namespace Libs::Graphics {
class CommandScheduler;
struct GraphicContext;
// The guest's query results (16-byte aligned): a begin/end pair of running z-pass counters per DB,
// 16 bytes apart; bit 63 marks a written counter. A query's samples are the sum of end - begin.
constexpr uint32_t OCCLUSION_QUERY_DBS = 16;
constexpr uint64_t OCCLUSION_QUERY_READY = 1ull << 63u;
// A dump of the running count, all of it in DB 0. A begin also clears its end counters.
inline void WriteZPassCounters(volatile uint64_t* results, uint64_t samples, bool begin) {
for (uint32_t db = 0; db < OCCLUSION_QUERY_DBS; db++) {
results[db * 2] = OCCLUSION_QUERY_READY | (db == 0 ? samples : 0);
if (begin) {
results[db * 2 + 1] = 0;
}
}
}
// The guest's z-pass counters from host occlusion queries: a draw counts its samples while a guest
// query is open, and the dumps' counters are written before the guest can see the GPU ran them.
class OcclusionQueries {
public:
OcclusionQueries(GraphicContext& graphics, CommandScheduler& scheduler);
~OcclusionQueries();
KYTY_CLASS_NO_COPY(OcclusionQueries);
[[nodiscard]] bool Enabled() const { return m_pool != nullptr; }
// In the draw's rendering: the query that counts its samples, or -1.
[[nodiscard]] int32_t Begin(vk::CommandBuffer command) {
return m_open == 0 ? -1 : BeginQuery(command);
}
void End(vk::CommandBuffer command, int32_t query) const {
if (query >= 0) {
command.endQuery(m_pool, static_cast<uint32_t>(query));
}
}
// A PIXEL_PIPE_STAT_DUMP to address, after the draws recorded so far.
void Dump(uint64_t address);
// Waits for the GPU to run the dumps and writes their counters.
void Publish() {
if (!m_dumps.empty()) {
PublishDumps();
}
}
private:
struct PendingDump {
uint64_t address;
uint64_t tick;
uint64_t queries;
uint64_t uncounted;
};
int32_t BeginQuery(vk::CommandBuffer command);
void PublishDumps();
CommandScheduler& m_scheduler;
vk::Device m_device;
vk::QueryPool m_pool;
vk::QueryControlFlags m_control;
std::vector<PendingDump> m_dumps;
std::vector<uint64_t> m_results;
uint32_t m_open = 0;
uint64_t m_begun = 0;
uint64_t m_read = 0;
uint64_t m_samples = 0;
uint64_t m_uncounted = 0;
};
} // namespace Libs::Graphics
#endif // EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_OCCLUSIONQUERIES_H_
@@ -585,6 +585,10 @@ TextureBinding RenderExecutor::ResolveTexture(const ShaderRecompiler::IR::ImageR
const auto alignment = desc.info.DataAlignment();
EXIT_NOT_IMPLEMENTED(desc.info.data.size == 0 || alignment == 0 ||
(address & (static_cast<uint64_t>(alignment) - 1u)) != 0);
// Image heaps keep the descriptors of unmapped textures, which PS5 never samples.
if (resource.heap && !m_context.IsMapped(address, desc.info.data.size)) {
return {.layout = vk::ImageLayout::eGeneral};
}
if (storage) {
ValidateStorageTexture(resource, descriptor, desc.info.data.size);
}
@@ -690,6 +694,10 @@ void RenderExecutor::PrepareBindings(const ShaderStageRuntime& runtime,
prepared.samplers.clear();
prepared.shader_data.clear();
for (uint32_t i = 0; i < program.info.images.size(); i++) {
if (program.info.images[i].heap) {
prepared.images[i] = {};
continue;
}
auto binding = ResolveTexture(program.info.images[i], snapshot.images[i]);
BindImage(binding.image_id, binding.desc.type == TextureCache::BindingType::Storage);
binding.mip_views.swap(prepared.images[i].mip_views);
@@ -852,14 +860,19 @@ void RenderExecutor::RebindImages(PreparedBindings& prepared) {
void RenderExecutor::PrepareGraphicsBindings(std::span<PreparedBindings* const> stages,
std::span<RenderColorInfo> colors) {
bool uses_dma = false;
bool uses_dma = false;
bool uses_image_heap = false;
FindBuffers(stages);
for (auto* stage: stages) {
uses_dma |= stage->runtime->program->info.uses_dma;
uses_image_heap |= stage->runtime->program->info.uses_image_heap;
}
if (uses_dma) {
m_context.PrepareBda();
}
if (uses_image_heap) {
m_context.GetImageHeap().Rebind(*this);
}
for (auto* stage: stages) {
RebindImages(*stage);
}
@@ -959,6 +972,9 @@ void RenderExecutor::CommitBindings(CommandBuffer& buffer,
}
}
}
if (pipeline.uses_image_heap) {
m_context.GetImageHeap().Commit(buffer, pipeline_bind_point, pipeline.pipeline_layout);
}
m_descriptor_images.resize(image_descriptor_count);
m_descriptor_writes.clear();
m_descriptor_writes.reserve(write_count);
@@ -1016,9 +1032,7 @@ void RenderExecutor::CommitBindings(CommandBuffer& buffer,
vk::AccessFlagBits2::eShaderRead | vk::AccessFlagBits2::eShaderWrite,
range, vk_buffer);
} else {
image.Transit(image.info.IsDepth() ? vk::ImageLayout::eDepthStencilReadOnlyOptimal
: vk::ImageLayout::eShaderReadOnlyOptimal,
vk::AccessFlagBits2::eShaderRead, range, vk_buffer);
image.Transit(image.ReadLayout(), vk::AccessFlagBits2::eShaderRead, range, vk_buffer);
}
binding.layout = image.backing.state.layout;
}
@@ -1089,6 +1103,7 @@ void RenderExecutor::CommitBindings(CommandBuffer& buffer,
for (uint32_t i = 0; i < program.info.images.size(); ++i) {
const auto& resource = program.info.images[i];
const auto& texture = descriptors.images[i];
if (resource.heap) continue;
EXIT_IF(texture.image_id &&
resource.mip_count != (texture.mip_views.empty() ? 1u : texture.mip_views.size()));
const auto kind = *ShaderRecompiler::IR::DescriptorBindingForImage(resource);
@@ -0,0 +1,544 @@
#include "graphics/host_gpu/renderer/pipeline/imageHeap.h"
#include "common/alignment.h"
#include "common/assert.h"
#include "common/profiler.h"
#include "graphics/host_gpu/pageManager.h"
#include "graphics/host_gpu/regionDefinitions.h"
#include "graphics/host_gpu/renderer/cache/streamBuffer.h"
#include "graphics/host_gpu/renderer/render.h"
#include "graphics/host_gpu/renderer/renderContext.h"
#include "graphics/shader/recompiler/ir/passes/ResourceMaterialization.h"
#include "kernel/memory.h"
#include <algorithm>
#include <xxhash.h>
namespace Libs::Graphics {
namespace {
namespace IR = ShaderRecompiler::IR;
// Smaller tables are cheap to materialize on each draw and are often transient.
constexpr uint64_t MinTableBytes = 64ull * 1024ull;
constexpr uint32_t ElementMask = (1u << IR::ImageHeapKindShift) - 1u;
template <typename Pred, typename Fn>
void ForEachRun(size_t begin, size_t end, Pred&& pred, Fn&& fn) {
while (begin < end) {
if (!pred(begin)) {
++begin;
continue;
}
auto run = begin + 1u;
while (run < end && pred(run)) ++run;
fn(begin, run);
begin = run;
}
}
// Grows [begin, end) to cover [first, last).
void Widen(size_t& begin, size_t& end, size_t first, size_t last) {
const bool empty = begin == end;
begin = empty ? first : std::min(begin, first);
end = empty ? last : std::max(end, last);
}
// Guest bytes of a table that its tracker pages [begin, end) cover.
std::pair<uint64_t, uint64_t> PageRange(uint64_t base, uint64_t size, size_t begin, size_t end) {
const auto origin = Common::AlignDown(base, TRACKER_PAGE_SIZE);
return {std::max(base, origin + begin * TRACKER_PAGE_SIZE),
std::min(base + size, origin + end * TRACKER_PAGE_SIZE)};
}
} // namespace
size_t ImageHeap::DescriptorHash::operator()(const DescriptorValue& value) const noexcept {
return XXH3_64bits(value.dwords.data(), sizeof(value.dwords));
}
ImageHeap::ImageHeap(RenderContext& context, PageManager& page_manager)
: m_context(context), m_page_manager(page_manager) {
auto& graphics = context.GetGraphics();
vk::PhysicalDeviceVulkan12Properties limits {};
vk::PhysicalDeviceProperties2 properties {};
properties.pNext = &limits;
graphics.physical_device.getProperties2(&properties);
EXIT_NOT_IMPLEMENTED(std::min({limits.maxDescriptorSetUpdateAfterBindSampledImages,
limits.maxPerStageDescriptorUpdateAfterBindSampledImages,
limits.maxPerStageUpdateAfterBindResources}) <
IR::ImageHeapCapacity);
const vk::DescriptorBindingFlags binding_flags =
vk::DescriptorBindingFlagBits::ePartiallyBound |
vk::DescriptorBindingFlagBits::eUpdateAfterBind |
vk::DescriptorBindingFlagBits::eUpdateUnusedWhilePending;
vk::DescriptorSetLayoutBindingFlagsCreateInfo flags_info {};
flags_info.bindingCount = 1;
flags_info.pBindingFlags = &binding_flags;
const vk::DescriptorSetLayoutBinding binding {0, vk::DescriptorType::eSampledImage,
IR::ImageHeapCapacity,
vk::ShaderStageFlagBits::eAll, nullptr};
vk::DescriptorSetLayoutCreateInfo layout_info {};
layout_info.pNext = &flags_info;
layout_info.flags = vk::DescriptorSetLayoutCreateFlagBits::eUpdateAfterBindPool;
layout_info.bindingCount = 1;
layout_info.pBindings = &binding;
EXIT_IF(graphics.device.createDescriptorSetLayout(&layout_info, nullptr, &m_layout) !=
vk::Result::eSuccess);
const vk::DescriptorPoolSize pool_size {vk::DescriptorType::eSampledImage,
IR::ImageHeapCapacity};
vk::DescriptorPoolCreateInfo pool_info {};
pool_info.flags = vk::DescriptorPoolCreateFlagBits::eUpdateAfterBind;
pool_info.maxSets = 1;
pool_info.poolSizeCount = 1;
pool_info.pPoolSizes = &pool_size;
EXIT_IF(graphics.device.createDescriptorPool(&pool_info, nullptr, &m_pool) !=
vk::Result::eSuccess);
vk::DescriptorSetAllocateInfo allocate {};
allocate.descriptorPool = m_pool;
allocate.descriptorSetCount = 1;
allocate.pSetLayouts = &m_layout;
EXIT_IF(graphics.device.allocateDescriptorSets(&allocate, &m_set) != vk::Result::eSuccess);
graphics.image_heap_layout = m_layout;
// Slot word 0 selects element 0, which stays a null image.
m_elements.push_back({.binding = {.layout = vk::ImageLayout::eShaderReadOnlyOptimal}});
m_writes.push_back(0u);
WriteDescriptors();
}
ImageHeap::~ImageHeap() {
for (const auto& memory: m_memories) Unwatch(*memory);
auto& graphics = m_context.GetGraphics();
graphics.image_heap_layout = nullptr;
graphics.device.destroyDescriptorPool(m_pool, nullptr);
graphics.device.destroyDescriptorSetLayout(m_layout, nullptr);
}
bool ImageHeap::Acquire(const IR::ImageHeapRequest& request, IR::ImageHeapView& result) {
KYTY_PROFILER_FUNCTION();
const auto size = request.size & ~uint64_t {3};
if (size < MinTableBytes || request.count == 0u || request.step == 0u ||
((request.base | request.first | request.step) & 3u) != 0u) {
return false;
}
auto memory = std::ranges::find_if(m_memories, [&](const auto& candidate) {
return candidate->base == request.base && candidate->size == size;
});
if (memory == m_memories.end()) {
if (!m_context.IsMapped(request.base, size)) return false;
auto created = std::make_unique<Memory>();
created->base = request.base;
created->size = size;
created->words.resize(size / 4u);
created->dead.resize(size / 4u);
created->dirty_pages.assign((Common::AlignUp(request.base + size, TRACKER_PAGE_SIZE) -
Common::AlignDown(request.base, TRACKER_PAGE_SIZE)) /
TRACKER_PAGE_SIZE,
uint8_t {1});
std::scoped_lock lock {m_mutex};
memory = m_memories.insert(m_memories.end(), std::move(created));
}
// A draw refreshes a table and a view takes its changes once, so all programs of the draw see
// one snapshot and one set of kinds.
if ((*memory)->draw != m_draw) {
if (!Refresh(**memory)) return false;
(*memory)->draw = m_draw;
}
auto& views = (*memory)->views;
const auto profile = IR::ImageHeapResource(*request.root, 0u);
auto view = std::ranges::find_if(views, [&](const auto& candidate) {
return candidate->step == request.step && candidate->first == request.first &&
candidate->count == request.count && candidate->profile == profile;
});
if (view == views.end()) {
const auto epoch = Image::heap_epoch.load(std::memory_order_relaxed);
auto created = std::make_unique<View>(View {.memory = **memory,
.step = request.step,
.first = request.first,
.count = request.count,
.profile = profile,
.resolve_epoch = epoch,
.layout_epoch = epoch});
created->slots.resize(size / 4u + 1u);
created->buffer = std::make_unique<Buffer>(
m_context.GetGraphics(), m_context.GetCommandScheduler(), MemoryUsage::DeviceLocal, 0,
vk::BufferUsageFlagBits::eTransferDst | vk::BufferUsageFlagBits::eStorageBuffer |
vk::BufferUsageFlagBits::eShaderDeviceAddress,
size + 4u);
UpdateProbes(*created, 0u, size);
// The first upload also clears the slots no probe selects.
created->dirty_begin = 0;
created->dirty_end = created->slots.size();
view = views.insert(views.end(), std::move(created));
} else {
// Empty after the first acquire in a draw, since the table refreshes once per draw.
UpdateProbes(**view, (*view)->changed_begin, (*view)->changed_end);
(*view)->changed_begin = (*view)->changed_end = 0;
}
if (std::ranges::find(m_active, view->get()) == m_active.end()) m_active.push_back(view->get());
const auto& selected = **view;
result = {.slots = selected.buffer->BufferDeviceAddress(),
.slot_count = static_cast<uint32_t>(selected.slots.size() - 1u),
.kinds = selected.kinds};
return true;
}
bool ImageHeap::Refresh(Memory& memory) {
// GPU writes leave no fault; reading the table downloads them.
const bool gpu_dirty = m_context.GetBufferCache().HasGpuDirtyBytes(memory.base, memory.size);
const auto pages = memory.dirty_pages.size();
m_runs.clear();
{
std::scoped_lock lock {m_mutex};
// Watch before reading, so a write during the read faults the page again.
ForEachRun(
0, pages, [&](size_t page) { return memory.dirty_pages[page] != 0u; },
[&](size_t begin, size_t end) {
std::fill(memory.dirty_pages.begin() + begin, memory.dirty_pages.begin() + end,
uint8_t {0});
Watch<true>(memory, begin, end);
m_runs.emplace_back(begin, end);
});
}
if (gpu_dirty) m_runs.assign(1u, {0u, pages});
for (size_t run = 0; run < m_runs.size(); ++run) {
const auto [address, end] = PageRange(memory.base, memory.size, m_runs[run].first,
m_runs[run].second);
m_scratch.resize((end - address) / 4u);
if (!LibKernel::Memory::TryReadBufferBacking(address, m_scratch.data(), end - address)) {
std::scoped_lock lock {m_mutex};
for (; run < m_runs.size(); ++run) {
MarkDirty(memory, m_runs[run].first, m_runs[run].second);
}
return false;
}
const auto words = std::span(memory.words).subspan((address - memory.base) / 4u,
m_scratch.size());
size_t first = 0;
size_t last = words.size();
while (first < last && words[first] == m_scratch[first]) ++first;
while (first < last && words[last - 1u] == m_scratch[last - 1u]) --last;
if (first == last) continue;
std::copy(m_scratch.begin() + first, m_scratch.begin() + last, words.begin() + first);
const auto offset = address - memory.base;
for (auto& view: memory.views) {
Widen(view->changed_begin, view->changed_end, offset + first * 4u, offset + last * 4u);
}
}
return true;
}
void ImageHeap::MarkDirty(Memory& memory, size_t begin, size_t end) {
ForEachRun(
begin, end, [&](size_t page) { return memory.dirty_pages[page] == 0u; },
[&](size_t first, size_t last) {
std::fill(memory.dirty_pages.begin() + first, memory.dirty_pages.begin() + last,
uint8_t {1});
Watch<false>(memory, first, last);
});
}
template <bool track>
void ImageHeap::Watch(const Memory& memory, size_t begin, size_t end) {
const auto [address, limit] = PageRange(memory.base, memory.size, begin, end);
m_page_manager.UpdatePageWatchers<track>(address, limit - address);
}
void ImageHeap::Unwatch(const Memory& memory) {
ForEachRun(
0, memory.dirty_pages.size(), [&](size_t page) { return memory.dirty_pages[page] == 0u; },
[&](size_t begin, size_t end) { Watch<false>(memory, begin, end); });
}
void ImageHeap::InvalidateMemory(uint64_t address, uint64_t size) {
std::scoped_lock lock {m_mutex};
for (auto& memory: m_memories) {
const auto begin = std::max(address, memory->base);
const auto end = std::min(address + size, memory->base + memory->size);
if (begin >= end) continue;
const auto origin = Common::AlignDown(memory->base, TRACKER_PAGE_SIZE);
MarkDirty(*memory, (begin - origin) / TRACKER_PAGE_SIZE,
(end - 1u - origin) / TRACKER_PAGE_SIZE + 1u);
}
}
void ImageHeap::UpdateProbes(View& view, uint64_t begin, uint64_t end) {
// A probe reads its descriptor from first + probe * step.
const uint64_t bytes = view.profile.r128 ? 16u : 32u;
auto probe = begin < view.first + bytes ? 0u : (begin - view.first - bytes) / view.step + 1u;
for (; probe < view.count && view.first + probe * view.step < end; ++probe) {
UpdateProbe(view, probe);
}
}
void ImageHeap::UpdateProbe(View& view, uint64_t probe) {
const auto& words = view.memory.words;
const auto offset = (view.first + probe * view.step) / 4u;
DescriptorValue value {.dword_count = 8u};
std::copy_n(words.begin() + offset,
std::min<size_t>(view.profile.r128 ? 4u : 8u, words.size() - offset),
value.dwords.begin());
// A buried word stays null in every view until its descriptor changes.
auto& dead = view.memory.dead[offset];
if (dead != 0u && dead != DescriptorHash {}(value)) dead = 0u;
const auto kind = dead != 0u ? 0u : IR::ImageHeapKind(view.profile, value);
Assign(view, offset, kind == 0u ? 0u : Intern(view, kind, value));
}
void ImageHeap::Assign(View& view, size_t slot, uint32_t entry) {
const auto previous = std::exchange(view.slots[slot], entry);
Release(view, previous);
if (previous != entry) Widen(view.dirty_begin, view.dirty_end, slot, slot + 1u);
}
void ImageHeap::Bury() {
if (m_unmapped.empty()) return;
for (const auto& memory: m_memories) {
for (const auto& view: memory->views) {
bool buried = false;
for (const auto& [value, index]: view->elements) {
const auto& binding = m_elements[index].binding;
if (binding.image_id && std::ranges::any_of(m_unmapped, [&](GuestRange range) {
return ImageRangeOverlaps(binding.desc.info.data, range);
})) {
m_elements[index].stale = buried = true;
}
}
if (!buried) continue;
for (uint64_t probe = 0; probe < view->count; ++probe) {
const auto slot = (view->first + probe * view->step) / 4u;
const auto& element = m_elements[view->slots[slot] & ElementMask];
if (!element.stale) continue;
memory->dead[slot] = DescriptorHash {}(element.value);
Assign(*view, slot, 0u);
}
}
}
m_unmapped.clear();
}
uint32_t ImageHeap::Intern(View& view, uint32_t kind, const DescriptorValue& value) {
const auto arm = static_cast<uint32_t>(std::ranges::find(view.kind_keys, kind) -
view.kind_keys.begin());
if (arm == view.kind_keys.size()) {
view.kind_keys.push_back(kind);
view.kinds.push_back(value);
}
auto [found, inserted] = view.elements.try_emplace(value, 0u);
if (inserted) {
found->second = Allocate();
auto& element = m_elements[found->second];
element.value = value;
element.kind = kind;
view.pending.push_back(found->second);
}
++m_elements[found->second].refs;
return arm << IR::ImageHeapKindShift | found->second;
}
void ImageHeap::Release(View& view, uint32_t entry) {
const auto index = entry & ElementMask;
if (index == 0u || --m_elements[index].refs != 0u) return;
auto& element = m_elements[index];
if (const auto found = view.elements.find(element.value);
found != view.elements.end() && found->second == index) {
view.elements.erase(found);
}
Unreference(element);
view.released.push_back(index);
}
void ImageHeap::Unreference(Element& element) {
if (auto* image = m_context.GetTextureCache().m_slot_images.try_get(element.binding.image_id)) {
image->heap_refs.fetch_sub(1u, std::memory_order_relaxed);
image->tick_accessed_last = m_context.GetCommandScheduler().CurrentTick();
}
}
uint32_t ImageHeap::Allocate() {
auto& scheduler = m_context.GetCommandScheduler();
while (!m_retired.empty() && scheduler.IsFree(m_retired.front().first)) {
m_free.push_back(m_retired.front().second);
m_retired.pop_front();
}
if (m_free.empty()) {
if (m_elements.size() == IR::ImageHeapCapacity) {
EXIT("image heaps exceed %u resident images\n", IR::ImageHeapCapacity);
}
m_elements.emplace_back();
return static_cast<uint32_t>(m_elements.size() - 1u);
}
const auto index = m_free.back();
m_free.pop_back();
m_elements[index] = {};
return index;
}
void ImageHeap::Retire(View& view) {
const auto tick = m_context.GetCommandScheduler().CurrentTick();
for (const auto index: view.released) m_retired.emplace_back(tick, index);
view.released.clear();
}
void ImageHeap::Rebind(RenderExecutor& executor) {
KYTY_PROFILER_FUNCTION();
auto& textures = m_context.GetTextureCache();
const auto epoch = Image::heap_epoch.load(std::memory_order_relaxed);
for (auto* view: m_active) {
bool stale = false;
const auto mark = [&](uint32_t index) { m_elements[index].stale = stale = true; };
if (view->resolve_epoch != epoch) {
view->resolve_epoch = epoch;
for (const auto& [value, index]: view->elements) {
const auto& binding = m_elements[index].binding;
if (!binding.image_id) continue;
const auto* image = textures.m_slot_images.try_get(binding.image_id);
if (image == nullptr || !image->registered || image->depth_id ||
image->binding.needs_rebind) {
mark(index);
} else {
(void)textures.FindTexture(binding.image_id, binding.desc);
}
}
}
// Discovery applies the color metadata clears written since the last draw.
std::erase_if(view->dcc, [&](uint32_t index) {
const auto& element = m_elements[index];
if (element.refs == 0u || element.stale) return true;
if (executor.ResolveTexture(IR::ImageHeapResource(view->profile, element.kind),
element.value).image_id == element.binding.image_id) {
return false;
}
mark(index);
return true;
});
if (stale) {
// A changed image needs a new element; pending draws still read the old one. The new
// element keeps the old value, since a view takes table changes only when acquired.
std::erase_if(view->elements,
[&](const auto& item) { return m_elements[item.second].stale; });
for (uint64_t probe = 0; probe < view->count; ++probe) {
const auto slot = (view->first + probe * view->step) / 4u;
const auto& element = m_elements[view->slots[slot] & ElementMask];
if (!element.stale) continue;
const auto value = element.value; // Intern may grow m_elements.
Assign(*view, slot, Intern(*view, element.kind, value));
}
}
for (const auto index: view->pending) {
if (m_elements[index].refs == 0u) continue;
auto binding = executor.ResolveTexture(
IR::ImageHeapResource(view->profile, m_elements[index].kind),
m_elements[index].value);
if (binding.image_id) {
auto& image = textures.GetImage(binding.image_id);
image.heap_refs.fetch_add(1u, std::memory_order_relaxed);
binding.layout = image.ReadLayout();
binding.image_view = textures.FindTexture(binding.image_id, binding.desc);
if (binding.desc.info.metadata.kind == ImageMetadataKind::Dcc) {
view->dcc.push_back(index);
}
view->away.push_back(index);
}
m_elements[index].binding = std::move(binding);
m_writes.push_back(index);
}
view->pending.clear();
}
WriteDescriptors();
}
void ImageHeap::WriteDescriptors() {
if (m_writes.empty()) return;
std::vector<vk::DescriptorImageInfo> infos;
std::vector<vk::WriteDescriptorSet> writes(m_writes.size());
infos.reserve(m_writes.size());
for (size_t write = 0; write < m_writes.size(); ++write) {
const auto& binding = m_elements[m_writes[write]].binding;
infos.push_back({nullptr, binding.image_view, binding.layout});
writes[write].dstSet = m_set;
writes[write].dstArrayElement = m_writes[write];
writes[write].descriptorCount = 1;
writes[write].descriptorType = vk::DescriptorType::eSampledImage;
writes[write].pImageInfo = &infos.back();
}
m_context.GetGraphics().device.updateDescriptorSets(static_cast<uint32_t>(writes.size()),
writes.data(), 0, nullptr);
m_writes.clear();
}
void ImageHeap::Commit(CommandBuffer& buffer, vk::PipelineBindPoint bind_point,
vk::PipelineLayout layout) {
KYTY_PROFILER_FUNCTION();
auto& textures = m_context.GetTextureCache();
const auto epoch = Image::heap_epoch.load(std::memory_order_relaxed);
for (auto* view: m_active) {
if (view->dirty_begin != view->dirty_end) {
const auto slots = std::span(view->slots).subspan(view->dirty_begin,
view->dirty_end - view->dirty_begin);
auto& stream = m_context.GetBufferCache().GetUtilityBuffer(MemoryUsage::Stream);
const auto offset = stream.Copy(slots.data(), slots.size_bytes(), sizeof(uint32_t));
view->buffer->CopyFrom(buffer, stream, offset, view->dirty_begin * sizeof(uint32_t),
slots.size_bytes());
view->dirty_begin = view->dirty_end = 0;
}
Retire(*view);
if (view->layout_epoch != epoch) {
view->layout_epoch = epoch;
view->away.clear();
for (const auto& [value, index]: view->elements) view->away.push_back(index);
}
// This draw's own bindings own the layouts of its attachments and storage images. An image
// this draw's discovery replaced lives until the tick ends, and its slots still select it.
std::erase_if(view->away, [&](uint32_t index) {
const auto& element = m_elements[index];
auto* image = textures.m_slot_images.try_get(element.binding.image_id);
if (element.refs == 0u || image == nullptr) return true;
if (image->binding.is_target || image->binding.shader_write ||
image->binding.force_general) {
return false;
}
const auto& view_info = element.binding.desc.view_info;
image->Transit(image->ReadLayout(), vk::AccessFlagBits2::eShaderRead,
ImageSubresourceRange {view_info.base_level, view_info.level_count,
view_info.base_layer, view_info.layer_count},
buffer.Handle());
return true;
});
}
m_active.clear();
buffer.Handle().bindDescriptorSets(bind_point, layout, IR::ImageHeapSet, 1, &m_set, 0,
nullptr);
}
void ImageHeap::UnmapMemory(uint64_t address, uint64_t size) {
// Descriptors of textures in the range stay null even once it is mapped again.
if (!m_unmapped.empty() && m_unmapped.back().End() == address) {
m_unmapped.back().size += size;
} else {
m_unmapped.push_back({address, size});
}
auto& scheduler = m_context.GetCommandScheduler();
std::scoped_lock lock {m_mutex};
std::erase_if(m_memories, [&](const std::unique_ptr<Memory>& memory) {
if (address >= memory->base + memory->size || memory->base >= address + size) {
return false;
}
Unwatch(*memory);
for (auto& view: memory->views) {
std::erase(m_active, view.get());
for (const auto& [value, index]: view->elements) {
Unreference(m_elements[index]);
view->released.push_back(index);
}
Retire(*view);
if (scheduler.Active()) {
scheduler.DeferOperation([buffer = std::move(view->buffer)] {});
}
}
return true;
});
}
} // namespace Libs::Graphics
@@ -0,0 +1,138 @@
#ifndef EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
#define EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
#include "common/common.h"
#include "graphics/host_gpu/renderer/pipeline/descriptors.h"
#include "graphics/shader/recompiler/ir/passes/SrtWalker.h"
#include <deque>
#include <memory>
#include <mutex>
#include <unordered_map>
#include <utility>
#include <vector>
namespace Libs::Graphics {
class Buffer;
class CommandBuffer;
class PageManager;
class RenderContext;
class RenderExecutor;
struct ImageHeapTestAccess;
// Guest image tables that shaders index by GPU-computed offsets, kept resident in one bindless
// sampled-image array. Each view of a table owns device slot words, one per table dword, that
// select an array element. CPU writes to a table fault in its snapshot, which each draw refreshes
// once; a view takes the changes when a draw first acquires it. Changes to the images behind
// elements bump Image::heap_epoch.
class ImageHeap final: public ShaderRecompiler::IR::ImageHeapSource {
public:
ImageHeap(RenderContext& context, PageManager& page_manager);
~ImageHeap();
KYTY_CLASS_NO_COPY(ImageHeap);
// Starts a draw or dispatch, before its programs acquire views.
void BeginDraw() {
++m_draw;
Bury();
}
bool Acquire(const ShaderRecompiler::IR::ImageHeapRequest& request,
ShaderRecompiler::IR::ImageHeapView& result) override;
// Image discovery: resolves new elements and rediscovers changed images.
void Rebind(RenderExecutor& executor);
// Recording: uploads slots, moves resident images to their read layout and binds the array.
void Commit(CommandBuffer& buffer, vk::PipelineBindPoint bind_point, vk::PipelineLayout layout);
// CPU write fault, on any thread.
void InvalidateMemory(uint64_t address, uint64_t size);
void UnmapMemory(uint64_t address, uint64_t size);
private:
using DescriptorValue = ShaderRecompiler::IR::DescriptorValue;
struct DescriptorHash {
size_t operator()(const DescriptorValue& value) const noexcept;
};
struct Element {
DescriptorValue value;
TextureBinding binding;
uint32_t kind = 0;
uint32_t refs = 0; // Slots of the owning view that select it.
bool stale = false;
};
struct Memory;
struct View {
Memory& memory;
uint64_t step = 0;
uint64_t first = 0;
uint64_t count = 0;
ShaderRecompiler::IR::ImageResource profile;
std::vector<uint32_t> slots; // Plus a trailing null slot.
std::vector<uint32_t> kind_keys;
std::vector<DescriptorValue> kinds; // One representative per kind_keys entry.
std::unordered_map<DescriptorValue, uint32_t, DescriptorHash> elements;
std::vector<uint32_t> pending; // Not yet resolved.
std::vector<uint32_t> released; // No slot selects them after the next upload.
std::vector<uint32_t> away; // Not yet moved to the read layout.
std::vector<uint32_t> dcc; // Color metadata clears apply on each discovery.
std::unique_ptr<Buffer> buffer;
size_t dirty_begin = 0; // Slots to upload.
size_t dirty_end = 0;
size_t changed_begin = 0; // Table bytes changed since it was probed.
size_t changed_end = 0;
uint64_t resolve_epoch = 0;
uint64_t layout_epoch = 0;
};
struct Memory {
uint64_t base = 0;
uint64_t size = 0;
std::vector<uint32_t> words;
std::vector<uint8_t> dirty_pages; // Unwatched, so unread since a write.
std::vector<uint64_t> dead; // Per word, a buried descriptor's hash.
uint64_t draw = 0; // Last draw that refreshed it.
std::vector<std::unique_ptr<View>> views;
};
[[nodiscard]] bool Refresh(Memory& memory);
void MarkDirty(Memory& memory, size_t begin, size_t end);
template <bool track>
void Watch(const Memory& memory, size_t begin, size_t end);
void Unwatch(const Memory& memory);
void UpdateProbes(View& view, uint64_t begin, uint64_t end);
void UpdateProbe(View& view, uint64_t probe);
void Assign(View& view, size_t slot, uint32_t entry);
void Bury();
uint32_t Intern(View& view, uint32_t kind, const DescriptorValue& value);
void Release(View& view, uint32_t entry);
void Unreference(Element& element);
uint32_t Allocate();
void Retire(View& view);
void WriteDescriptors();
RenderContext& m_context;
PageManager& m_page_manager;
std::mutex m_mutex; // Fault-visible Memory state and m_memories.
std::vector<std::unique_ptr<Memory>> m_memories;
std::vector<View*> m_active; // Acquired since the last commit.
uint64_t m_draw = 1;
std::vector<Element> m_elements;
std::vector<uint32_t> m_free;
std::deque<std::pair<uint64_t, uint32_t>> m_retired; // Tick, element.
std::vector<std::pair<size_t, size_t>> m_runs;
std::vector<uint32_t> m_scratch;
std::vector<uint32_t> m_writes; // Elements whose descriptors changed.
std::vector<GuestRange> m_unmapped; // Since the last Bury.
vk::DescriptorSetLayout m_layout = nullptr;
vk::DescriptorPool m_pool = nullptr;
vk::DescriptorSet m_set = nullptr;
friend struct ImageHeapTestAccess;
};
} // namespace Libs::Graphics
#endif // EMULATOR_SRC_GRAPHICS_HOST_GPU_RENDERER_IMAGEHEAP_H_
@@ -353,6 +353,7 @@ struct PipelineCache::ProgramCache {
.user_data = user_data,
.shader_base = params.Base(),
.read_specialization_memory = ReadShaderGuestMemory,
.image_heap = image_heap,
};
if constexpr (std::is_same_v<InputInfo, ShaderComputeInputInfo>) {
runtime.workgroup_counts = input_info.workgroup_counts;
@@ -480,13 +481,14 @@ struct PipelineCache::ProgramCache {
return permutation.handle;
}
explicit ProgramCache(GraphicContext& graphics)
: graphics(graphics),
ProgramCache(GraphicContext& graphics, ShaderRecompiler::IR::ImageHeapSource* image_heap)
: graphics(graphics), image_heap(image_heap),
// AMD GPUs (vendor 0x1002) run NMin/NMax as v_min/v_max, which give the guest's NaN and
// signed-zero results; GLSL.std.450 alone does not order -0 below +0.
profile {.maintenance8 = graphics.maintenance8_enabled,
.native_float_min_max =
graphics.physical_device_properties.vendorID == 0x1002u} {
graphics.physical_device_properties.vendorID == 0x1002u,
.device_clock_period_ps = graphics.device_clock_period_ps} {
lookup_key.static_state.reserve(MaxStaticKeyWords);
}
~ProgramCache() {
@@ -501,12 +503,14 @@ struct PipelineCache::ProgramCache {
std::unordered_map<ProgramKey, SourceEntry, ProgramKeyHash> programs;
ProgramKey lookup_key;
GraphicContext& graphics;
ShaderRecompiler::IR::ImageHeapSource* image_heap;
const ShaderRecompiler::Spirv::Profile profile;
uint64_t next_shader_id = 0;
};
PipelineCache::PipelineCache(GraphicContext& graphics)
: m_graphics(graphics), m_program_cache(std::make_unique<ProgramCache>(graphics)),
PipelineCache::PipelineCache(GraphicContext& graphics,
ShaderRecompiler::IR::ImageHeapSource* image_heap)
: m_graphics(graphics), m_program_cache(std::make_unique<ProgramCache>(graphics, image_heap)),
m_next_driver_cache_save(std::chrono::steady_clock::now() + DriverCacheSaveInterval) {
EXIT_NOT_IMPLEMENTED(!Common::Thread::IsMainThread());
InitializeDriverCache();
@@ -30,6 +30,10 @@ class UserConfig;
struct ComputeShaderInfo;
} // namespace HW
namespace ShaderRecompiler::IR {
class ImageHeapSource;
} // namespace ShaderRecompiler::IR
#pragma pack(push, 1)
struct PipelineStaticParameters {
@@ -104,7 +108,7 @@ struct ShaderProgram {
// The owning renderer serializes access, including saves while the GPU is running.
class PipelineCache {
public:
explicit PipelineCache(GraphicContext& graphics);
PipelineCache(GraphicContext& graphics, ShaderRecompiler::IR::ImageHeapSource* image_heap);
~PipelineCache();
KYTY_CLASS_NO_COPY(PipelineCache);
void Save();
@@ -114,6 +118,7 @@ public:
vk::Pipeline pipeline = nullptr;
vk::DescriptorSetLayout descriptor_set_layout = nullptr;
bool uses_push_descriptors = false;
bool uses_image_heap = false;
};
struct GraphicsPrograms {
@@ -473,20 +473,23 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
for (const auto& stage: vertex_info) {
const auto native_stage = NativeShaderStage(stage.logical_stage);
AddLayoutBindings(descriptor_bindings, *stage.stage.program, native_stage);
pipeline.uses_image_heap |= stage.stage.program->info.uses_image_heap;
graphics_stages |= native_stage;
}
if (ps_active) {
EXIT_IF(!ps_input_info->stage);
AddLayoutBindings(descriptor_bindings, *ps_input_info->stage.program,
vk::ShaderStageFlagBits::eFragment);
pipeline.uses_image_heap |= ps_input_info->stage.program->info.uses_image_heap;
}
CreateDescriptorLayout(graphics, pipeline, descriptor_bindings);
const vk::PushConstantRange push_constants {graphics_stages, 0,
ShaderRecompiler::IR::NativePushConstantSize};
const std::array set_layouts {pipeline.descriptor_set_layout, graphics.image_heap_layout};
vk::PipelineLayoutCreateInfo pipeline_layout_info {};
pipeline_layout_info.setLayoutCount = 1;
pipeline_layout_info.pSetLayouts = &pipeline.descriptor_set_layout;
pipeline_layout_info.setLayoutCount = pipeline.uses_image_heap ? 2u : 1u;
pipeline_layout_info.pSetLayouts = set_layouts.data();
pipeline_layout_info.pushConstantRangeCount = 1;
pipeline_layout_info.pPushConstantRanges = &push_constants;
@@ -615,13 +618,15 @@ void CreatePipelineInternal(GraphicContext& graphics, PipelineCache::Pipeline& p
std::vector<vk::DescriptorSetLayoutBinding> descriptor_bindings;
AddLayoutBindings(descriptor_bindings, *input_info.stage.program,
vk::ShaderStageFlagBits::eCompute);
pipeline.uses_image_heap = input_info.stage.program->info.uses_image_heap;
CreateDescriptorLayout(graphics, pipeline, descriptor_bindings);
const vk::PushConstantRange push_constants {vk::ShaderStageFlagBits::eCompute, 0,
ShaderRecompiler::IR::NativePushConstantSize};
const std::array set_layouts {pipeline.descriptor_set_layout, graphics.image_heap_layout};
vk::PipelineLayoutCreateInfo pipeline_layout_info {};
pipeline_layout_info.setLayoutCount = 1;
pipeline_layout_info.pSetLayouts = &pipeline.descriptor_set_layout;
pipeline_layout_info.setLayoutCount = pipeline.uses_image_heap ? 2u : 1u;
pipeline_layout_info.pSetLayouts = set_layouts.data();
pipeline_layout_info.pushConstantRangeCount = 1;
pipeline_layout_info.pPushConstantRanges = &push_constants;
+2 -2
View File
@@ -171,6 +171,8 @@ public:
void CommitBindings(CommandBuffer& buffer, vk::PipelineBindPoint pipeline_bind_point,
const PipelineCache::Pipeline& pipeline,
std::span<PreparedBindings* const> bindings);
[[nodiscard]] TextureBinding ResolveTexture(const ShaderRecompiler::IR::ImageResource& resource,
const ShaderRecompiler::IR::DescriptorValue& value);
private:
// False, with nothing drawn, when the draw's indirect records need CPU counts.
@@ -182,8 +184,6 @@ private:
std::optional<PreparedBindings> pixel;
};
[[nodiscard]] TextureBinding ResolveTexture(const ShaderRecompiler::IR::ImageResource& resource,
const ShaderRecompiler::IR::DescriptorValue& value);
void PrepareGraphicsBindings(std::span<PreparedBindings* const> stages,
std::span<RenderColorInfo> colors);
void ResolveRenderColorTarget(CommandBuffer& buffer, RenderColorInfo& target,
@@ -311,6 +311,7 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
input_info.workgroup_counts[axis] = threads / size + (threads % size != 0u);
}
}
m_context.GetImageHeap().BeginDraw();
const auto compute_program =
m_context.GetPipelineCache().GetComputeProgram(cs_regs, sh_regs, input_info);
if (!compute_program) return;
@@ -363,6 +364,9 @@ void RenderExecutor::DispatchDirect(uint64_t submit_id, CommandBuffer& buffer,
if (program.info.uses_dma) {
m_context.PrepareBda();
}
if (program.info.uses_image_heap) {
m_context.GetImageHeap().Rebind(*this);
}
RebindImages(bindings);
BindSharedMemory(m_context, input_info, bindings);
RebindBuffers(bindings);
@@ -497,6 +501,7 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
ShaderComputeInputInfo input_info {};
input_info.async_compute = async_compute;
input_info.indirect_args_address = args_addr;
m_context.GetImageHeap().BeginDraw();
const auto compute_program = m_context.GetPipelineCache().GetComputeProgram(
cs_regs, buffer.GetRegisters().GetShaderRegisters(), input_info);
if (!compute_program) return;
@@ -511,6 +516,9 @@ void RenderExecutor::DispatchIndirect(uint64_t submit_id, CommandBuffer& buffer,
m_context.PrepareBda();
}
BindSharedMemory(m_context, input_info, bindings, args_addr);
if (program.info.uses_image_heap) {
m_context.GetImageHeap().Rebind(*this);
}
RebindImages(bindings);
// Acquiring arguments can merge cache buffers; finalize shader bindings afterward.
const auto [args_buffer, args_offset] = m_context.GetBufferCache().ObtainBuffer(
@@ -13,9 +13,10 @@ namespace Libs::Graphics {
RenderContext::RenderContext(GraphicContext& graphics)
: m_graphics(graphics), m_render_executor(*this), m_command_scheduler(*this, graphics),
m_descriptor_heap(graphics, m_command_scheduler.GetMasterSemaphore()),
m_pipeline_cache(graphics), m_sampler_cache(graphics),
m_pipeline_cache(graphics, &m_image_heap), m_sampler_cache(graphics),
m_buffer_cache(graphics, m_command_scheduler, m_page_manager, m_texture_cache),
m_texture_cache(graphics, m_command_scheduler, m_page_manager, m_buffer_cache) {
m_texture_cache(graphics, m_command_scheduler, m_page_manager, m_buffer_cache),
m_image_heap(*this, m_page_manager), m_occlusion_queries(graphics, m_command_scheduler) {
EXIT_NOT_IMPLEMENTED(!Common::Thread::IsMainThread());
}
@@ -64,6 +65,7 @@ bool RenderContext::HandleFault(PageFaultAccess access, uint64_t fault_vaddr) no
if (access == PageFaultAccess::Write) {
m_buffer_cache.InvalidateMemory(fault_vaddr, fault_size);
m_texture_cache.InvalidateMemory(fault_vaddr, fault_size);
m_image_heap.InvalidateMemory(fault_vaddr, fault_size);
} else {
m_buffer_cache.ReadMemory(fault_vaddr, fault_size);
}
@@ -76,6 +78,7 @@ bool RenderContext::InvalidateMemory(uint64_t vaddr, uint64_t size) {
}
m_buffer_cache.InvalidateMemory(vaddr, size);
m_texture_cache.InvalidateMemory(vaddr, size);
m_image_heap.InvalidateMemory(vaddr, size);
return true;
}
@@ -112,6 +115,7 @@ void RenderContext::UnmapMemory(uint64_t vaddr, uint64_t size) {
}
m_buffer_cache.InvalidateMemory(vaddr, size);
m_texture_cache.UnmapMemory(vaddr, size);
m_image_heap.UnmapMemory(vaddr, size);
std::lock_guard lock(m_mapped_ranges_mutex);
m_mapped_ranges.Subtract(vaddr, size);
};
@@ -11,7 +11,9 @@
#include "graphics/host_gpu/renderer/cache/samplerCache.h"
#include "graphics/host_gpu/renderer/cache/textureCache.h"
#include "graphics/host_gpu/renderer/commandScheduler.h"
#include "graphics/host_gpu/renderer/occlusionQueries.h"
#include "graphics/host_gpu/renderer/pipeline/descriptorHeap.h"
#include "graphics/host_gpu/renderer/pipeline/imageHeap.h"
#include "graphics/host_gpu/renderer/pipeline/pipelineCache.h"
#include "kernel/eventQueue.h"
@@ -43,10 +45,12 @@ public:
CommandScheduler& GetCommandScheduler() { return m_command_scheduler; }
PipelineCache& GetPipelineCache() { return m_pipeline_cache; }
DescriptorHeap& GetDescriptorHeap() { return m_descriptor_heap; }
ImageHeap& GetImageHeap() { return m_image_heap; }
SamplerCache& GetSamplerCache() { return m_sampler_cache; }
BufferCache& GetBufferCache() { return m_buffer_cache; }
TextureCache& GetTextureCache() { return m_texture_cache; }
RenderExecutor& GetRenderExecutor() { return m_render_executor; }
OcclusionQueries& GetOcclusionQueries() { return m_occlusion_queries; }
[[nodiscard]] bool HandleFault(PageFaultAccess access, uint64_t fault_vaddr) noexcept;
[[nodiscard]] bool InvalidateMemory(uint64_t vaddr, uint64_t size);
@@ -76,6 +80,8 @@ private:
PageManager m_page_manager;
BufferCache m_buffer_cache;
TextureCache m_texture_cache;
ImageHeap m_image_heap;
OcclusionQueries m_occlusion_queries;
mutable std::shared_mutex m_mapped_ranges_mutex;
RangeSet m_mapped_ranges;
std::unique_ptr<GuestGpu> m_gpu;
@@ -893,6 +893,7 @@ static void RefreshShaders(CommandBuffer& buffer, const DrawCallInfo& draw,
}
}
auto& pipeline_cache = buffer.GetContext().GetPipelineCache();
buffer.GetContext().GetImageHeap().BeginDraw();
if (draw.IsIndexed()) {
LogDrawPhase(draw.Name(), "GetGraphicsPrograms");
}
@@ -1283,6 +1284,8 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
if (!draw.IsIndexed()) {
SetDrawDebugPhase(buffer, submit_id, draw, 0x500u);
}
auto& occlusion_queries = m_context.GetOcclusionQueries();
const int32_t occlusion_query = occlusion_queries.Begin(vk_buffer);
if (mesh_active) {
vk_buffer.drawMeshTasksEXT(mesh_groups, draw.instance_count, 1);
} else if (args_buffer.first == nullptr) {
@@ -1294,6 +1297,7 @@ void RenderExecutor::ExecutePreparedDraw(uint64_t submit_id, CommandBuffer& buff
vk_buffer.drawIndirect(args_buffer.first->Handle(), args_buffer.second,
emit.indirect.count, emit.indirect.stride);
}
occlusion_queries.End(vk_buffer, occlusion_query);
Common::DebugCounters::Add(Common::DebugCounters::Counter::Draws);
Common::DebugCounters::Add(Common::DebugCounters::Counter::DrawInstances, draw.instance_count);
@@ -24,6 +24,7 @@
#include <SDL3/SDL.h>
#include <SDL3/SDL_vulkan.h>
#include <algorithm>
#include <cmath>
#include <cstdio>
#include <cstring>
#include <fmt/format.h>
@@ -65,6 +66,9 @@ vk::PhysicalDeviceVulkan12Features WindowContext::RequiredVulkan12Features() noe
features.shaderSampledImageArrayNonUniformIndexing = VK_TRUE;
features.shaderStorageImageArrayNonUniformIndexing = VK_TRUE;
features.shaderStorageBufferArrayNonUniformIndexing = VK_TRUE;
features.descriptorBindingSampledImageUpdateAfterBind = VK_TRUE;
features.descriptorBindingUpdateUnusedWhilePending = VK_TRUE;
features.descriptorBindingPartiallyBound = VK_TRUE;
return features;
}
@@ -295,6 +299,14 @@ static void VulkanFindPhysicalDevice(vk::Instance instance, vk::SurfaceKHR surfa
check_feature(features12.shaderStorageBufferArrayNonUniformIndexing,
"shaderStorageBufferArrayNonUniformIndexing",
required_features12.shaderStorageBufferArrayNonUniformIndexing);
check_feature(features12.descriptorBindingSampledImageUpdateAfterBind,
"descriptorBindingSampledImageUpdateAfterBind",
required_features12.descriptorBindingSampledImageUpdateAfterBind);
check_feature(features12.descriptorBindingUpdateUnusedWhilePending,
"descriptorBindingUpdateUnusedWhilePending",
required_features12.descriptorBindingUpdateUnusedWhilePending);
check_feature(features12.descriptorBindingPartiallyBound, "descriptorBindingPartiallyBound",
required_features12.descriptorBindingPartiallyBound);
check_feature(features13.robustImageAccess, "robustImageAccess");
check_feature(robustness2.nullDescriptor, "nullDescriptor");
check_feature(features13.dynamicRendering, "dynamicRendering",
@@ -495,6 +507,13 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
maintenance8.pNext = supported_features2.pNext;
supported_features2.pNext = &maintenance8;
}
const bool shader_clock_extension =
HasExtension(device_extensions, VK_KHR_SHADER_CLOCK_EXTENSION_NAME);
vk::PhysicalDeviceShaderClockFeaturesKHR shader_clock {};
if (shader_clock_extension) {
shader_clock.pNext = supported_features2.pNext;
supported_features2.pNext = &shader_clock;
}
const bool image_atomic_int64_extension =
HasExtension(device_extensions, VK_EXT_SHADER_IMAGE_ATOMIC_INT64_EXTENSION_NAME);
vk::PhysicalDeviceShaderImageAtomicInt64FeaturesEXT image_atomic_int64 {};
@@ -582,6 +601,11 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
feedback_extensions && feedback_layout.attachmentFeedbackLoopLayout &&
feedback_dynamic.attachmentFeedbackLoopDynamicState;
graphics.maintenance8_enabled = maintenance8_extension && maintenance8.maintenance8;
// The device clock ticks at timestampPeriod nanoseconds.
if (shader_clock_extension && shader_clock.shaderDeviceClock) {
graphics.device_clock_period_ps = static_cast<uint32_t>(
std::lround(graphics.physical_device_properties.limits.timestampPeriod * 1000.0f));
}
graphics.conditional_rendering_enabled =
conditional_rendering_extension && conditional_rendering.conditionalRendering;
LOGF("Vulkan maintenance8 support: %s\n",
@@ -631,6 +655,16 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
vk::SampleCountFlagBits::e1;
}
graphics.sample_rate_shading_enabled = true;
// Occlusion queries reset their queries from the host and count samples exactly where they can.
graphics.occlusion_queries_enabled =
Config::OcclusionQueriesEnabled() && supported_features12.hostQueryReset;
graphics.occlusion_query_precise_enabled =
graphics.occlusion_queries_enabled && supported_features2.features.occlusionQueryPrecise;
features12.hostQueryReset = graphics.occlusion_queries_enabled;
device_features.occlusionQueryPrecise = graphics.occlusion_query_precise_enabled;
LOGF("Vulkan occlusion queries: %s, precise: %s\n",
graphics.occlusion_queries_enabled ? "true" : "false",
graphics.occlusion_query_precise_enabled ? "true" : "false");
device_features.shaderInt64 = VK_TRUE;
device_features.shaderFloat64 =
supported_features2.features.shaderFloat64 &&
@@ -691,6 +725,10 @@ static vk::Device VulkanCreateDevice(GraphicContext& graphics,
maintenance8.pNext = const_cast<void*>(create_info.pNext);
create_info.pNext = &maintenance8;
}
if (graphics.device_clock_period_ps != 0u) {
shader_clock.pNext = const_cast<void*>(create_info.pNext);
create_info.pNext = &shader_clock;
}
if (graphics.conditional_rendering_enabled) {
conditional_rendering.pNext = const_cast<void*>(create_info.pNext);
create_info.pNext = &conditional_rendering;
@@ -1066,6 +1104,7 @@ void WindowContext::CreateVulkan() {
graphic_ctx.memory_budget_ext_enabled = true;
}
for (const auto* extension: {VK_KHR_MAINTENANCE_8_EXTENSION_NAME,
VK_KHR_SHADER_CLOCK_EXTENSION_NAME,
VK_EXT_SHADER_IMAGE_ATOMIC_INT64_EXTENSION_NAME,
VK_EXT_PROVOKING_VERTEX_EXTENSION_NAME,
VK_EXT_MESH_SHADER_EXTENSION_NAME,
@@ -32,8 +32,6 @@
#include <vector>
#include <vulkan/vk_platform.h>
#define STB_IMAGE_IMPLEMENTATION
#define STBI_NO_SIMD
#include "stb_image.h"
// IWYU pragma: no_include <intrin.h>
@@ -161,6 +161,7 @@ void ValidateNativeProgram(const IR::Program& program, bool lds_storage) {
if (!kind.has_value()) {
Fail(program, "native shader plan has an invalid image class");
}
if (image.heap) continue;
if (image.mip_count == 0u ||
(image.mip_mode != IR::ImageMipMode::Dynamic && image.mip_count != 1u)) {
Fail(program, "native shader plan has an invalid image mip descriptor count");
@@ -13,6 +13,8 @@ namespace Libs::Graphics::ShaderRecompiler::Spirv {
struct Profile {
bool maintenance8 = false; // Dynamic texel offsets on image samples.
bool native_float_min_max = false; // NMin/NMax give the guest's NaN and signed-zero results.
// OpReadClockKHR Device tick in picoseconds; 0 without shaderDeviceClock.
uint32_t device_clock_period_ps = 0;
};
inline constexpr Profile DefaultProfile {};
@@ -120,6 +120,10 @@ uint32_t ImageType(EmitterState& state, const IR::ImageResource& image) {
} else {
EXIT("invalid image resource class");
}
if (image.dimension == ImageDimension::Dim1D || image.dimension == ImageDimension::Dim1DArray) {
state.builder.RequireCapability(sampled == 1 ? spv::CapabilitySampled1D
: spv::CapabilityImage1D);
}
const auto& info = ImageDimensionInfoFor(image.dimension);
if (sampled == 2 && info.multisampled != 0u) {
state.builder.RequireCapability(spv::CapabilityStorageImageMultisample);
@@ -147,14 +151,20 @@ void DecorateNonUniformDescriptor(EmitterState& state, uint32_t id, spv::Capabil
state.builder.RequireCapability(capability);
state.builder.AddAnnotation(spv::OpDecorate, id, spv::DecorationNonUniform);
}
ImageDefinition& ImageDescriptorDefinition(EmitterState& state, const IR::ImageResource& image,
IR::DescriptorBindingKind kind) {
return (image.heap ? state.heap_images : state.images)[IR::ImageBindingIndex(kind)];
}
} // namespace
uint32_t LoadImageDescriptor(EmitterState& state, uint32_t resource, uint32_t mip,
uint32_t array_index) {
const auto pointer = ImageDescriptorPointer(state, resource, mip, array_index);
const auto image = state.builder.AllocateId();
const auto kind = *IR::DescriptorBindingForImage(state.program.info.images[resource]);
state.builder.AddFunction(spv::OpLoad, state.images[IR::ImageBindingIndex(kind)].type,
const auto& resource_info = state.program.info.images[resource];
const auto kind = *IR::DescriptorBindingForImage(resource_info);
state.builder.AddFunction(spv::OpLoad, ImageDescriptorDefinition(state, resource_info, kind).type,
image, pointer);
if (array_index != 0u) {
DecorateNonUniformDescriptor(state, image,
@@ -187,13 +197,13 @@ uint32_t MakeSampledImage(EmitterState& state, uint32_t resource, uint32_t sampl
EXIT_IF(image_resource.resource_class != IR::ImageResourceClass::Sampled);
const auto image = LoadImageDescriptor(state, resource, mip, array_index);
const auto sampled_image = state.builder.AllocateId();
const auto kind = *IR::DescriptorBindingForImage(image_resource);
auto& sampled_type = state.images[IR::ImageBindingIndex(kind)].sampled_type;
if (sampled_type == 0) {
sampled_type = state.builder.Type(
spv::OpTypeSampledImage, state.images[IR::ImageBindingIndex(kind)].type);
auto& definition = ImageDescriptorDefinition(state, image_resource,
*IR::DescriptorBindingForImage(image_resource));
if (definition.sampled_type == 0) {
definition.sampled_type = state.builder.Type(spv::OpTypeSampledImage, definition.type);
}
state.builder.AddFunction(spv::OpSampledImage, sampled_type, sampled_image, image, sampler_id);
state.builder.AddFunction(spv::OpSampledImage, definition.sampled_type, sampled_image, image,
sampler_id);
if (array_index != 0u || sampler_dynamic) {
DecorateNonUniformDescriptor(state, sampled_image, spv::CapabilitySampledImageArrayNonUniformIndexing);
}
@@ -205,11 +215,11 @@ uint32_t ImageDescriptorPointer(EmitterState& state, uint32_t resource, uint32_t
const auto& image = state.program.info.images.at(resource);
EXIT_IF(mip >= image.mip_count);
const auto kind = IR::DescriptorBindingForImage(image);
EXIT_IF(!kind.has_value());
EXIT_IF(!kind.has_value() || (image.heap && array_index == 0u));
if (array_index == 0u) {
array_index = ConstantU32(state, image.descriptor_index + mip);
}
const auto& definition = state.images[IR::ImageBindingIndex(*kind)];
const auto& definition = ImageDescriptorDefinition(state, image, *kind);
return DescriptorElementPointer(state, definition.pointer_type, definition.variable,
array_index, *kind, resource,
"image descriptor array was not emitted");
@@ -4,6 +4,7 @@
#include <algorithm>
#include <atomic>
#include <bit>
#include <optional>
namespace Libs::Graphics::ShaderRecompiler::Spirv::Emitter {
@@ -132,7 +133,7 @@ uint32_t EmitDppWriteCondition(ValueEmitContext& ctx, const IR::DppMoveFlags& fl
enable(flags.bank_mask, 2u);
enable(flags.row_mask, 4u);
// Only row shifts (0x101-0x10f, 0x111-0x11f) can vacate a lane; other controls always read one.
const bool row_shift = !flags.dpp8 && flags.control > 0x100u && flags.control < 0x120u &&
const bool row_shift = flags.control > 0x100u && flags.control < 0x120u &&
flags.control != 0x110u;
if (!flags.bound_control && row_shift) {
writable = Binary(state, spv::OpLogicalAnd, TypeBool(state), writable,
@@ -881,6 +882,22 @@ uint32_t EmitGetShaderBase(ValueEmitContext& ctx) {
return ctx.Def(IR::Value(uint64_t {0}));
}
uint32_t EmitReadRealtimeClock(ValueEmitContext& ctx, const IR::Inst& inst) {
auto& state = ctx.state;
const auto period = state.profile.device_clock_period_ps;
if (period == 0u) {
ctx.Fail(inst, "requires shaderDeviceClock");
}
state.builder.RequireExtension("SPV_KHR_shader_clock");
state.builder.RequireCapability(spv::CapabilityShaderClockKHR);
const auto clock =
Unary(state, spv::OpReadClockKHR, TypeU64(state), ConstantU32(state, spv::ScopeDevice));
// S_MEMREALTIME counts at 100 MHz. Shift the host clock by the most bits that keep its tick at
// most 15000 ps (2/3 of 100 MHz): 1 GHz shifts by 3 to 125 MHz; 100 MHz and slower stay.
const uint32_t shift = std::bit_width(7500u / period);
return EmitShiftRightLogical64(ctx, clock, IR::Value(shift));
}
void EmitUnreachable(ValueEmitContext& ctx, const IR::Inst& inst) {
ctx.Fail(inst, "must be lowered before SPIR-V emission");
}
@@ -700,6 +700,21 @@ auto EmitImageAccess(ValueEmitContext& ctx, const IR::Inst& inst, uint32_t resul
image.indirect_resources.size() < 2u) {
ctx.Fail(inst, "has no indirect image runtime mapping");
}
if (image.heap) {
const auto entry = EmitImageHeapEntry(state, key, image.indirect_mapping_offset);
const auto kinds = static_cast<uint32_t>(image.indirect_resources.size()) - 1u;
// One kind leaves every slot word a bare element.
if (kinds == 1u) return emit(image.indirect_resources[1], entry);
const auto element = Binary(state, spv::OpBitwiseAnd, TypeU32(state), entry,
ConstantU32(state, (1u << IR::ImageHeapKindShift) - 1u));
const auto EmitKind = [&](uint32_t kind) {
return emit(image.indirect_resources[kind + 1u], element);
};
return EmitIndexSwitch(state,
Binary(state, spv::OpShiftRightLogical, TypeU32(state), entry,
ConstantU32(state, IR::ImageHeapKindShift)),
kinds, result_type, EmitKind);
}
const auto selected = EmitIndirectResourceIndex(
state, key, image.indirect_mapping_offset, image.indirect_search_iterations, 0u);
const auto SameType = [&](uint32_t lhs, uint32_t rhs) {
@@ -879,14 +894,19 @@ void EmitImage(ValueEmitContext& ctx, const IR::Inst& inst) {
}
if (op == IR::ValueOpcode::ImageQueryLod) {
state.builder.RequireCapability(spv::CapabilityImageQuery);
const auto dimension = image.dimension;
const auto sampled = MakeSampledImage(state, mem.resource,
LoadSamplerDescriptor(state, mem.sampler));
const auto lod = state.builder.AllocateId();
state.builder.AddFunction(
spv::OpImageQueryLod, TypeF32Vector(state, 2), lod, sampled,
CoordF32(ctx, mem, *address, 0, ImageDimensionInfoFor(dimension).spatial_components,
image.cube));
const auto sampler = LoadSamplerDescriptor(state, mem.sampler);
const auto lod = EmitImageAccess(ctx, inst, TypeF32Vector(state, 2),
[&](uint32_t resource, uint32_t array_index) {
const auto& candidate = state.program.info.images[resource];
const auto query = state.builder.AllocateId();
state.builder.AddFunction(
spv::OpImageQueryLod, TypeF32Vector(state, 2), query,
MakeSampledImage(state, resource, sampler, 0u, array_index),
CoordF32(ctx, mem, *address, 0,
ImageDimensionInfoFor(candidate.dimension).spatial_components,
candidate.cube));
return query;
});
uint32_t values[4] = {ConstantU32(state, 0), ConstantU32(state, 0), ConstantU32(state, 0),
ConstantU32(state, 0)};
for (uint32_t index = 0; index < 2u; index++) {
@@ -240,6 +240,7 @@ uint32_t EmitGetAttributeF16(ValueEmitContext& ctx, const IR::Inst&
uint32_t EmitGetInterpolationParameter(ValueEmitContext& ctx, const IR::Inst& inst);
void EmitSetAttribute(ValueEmitContext& ctx, const IR::Inst& inst);
uint32_t EmitGetShaderBase(ValueEmitContext& ctx);
uint32_t EmitReadRealtimeClock(ValueEmitContext& ctx, const IR::Inst& inst);
inline constexpr auto EmitGetSrtResource = EmitVoid;
inline constexpr auto EmitGetBufferResource = EmitGetSrtResource;
uint32_t EmitGetBufferResourceIndex(ValueEmitContext& ctx, const IR::Inst& inst);
@@ -130,6 +130,7 @@ struct EmitterState {
uint32_t lds_u64_variable = 0;
std::array<uint32_t, 2> scratch_variable {};
std::array<ImageDefinition, IR::ImageBindingCount> images {};
std::array<ImageDefinition, IR::ImageBindingCount> heap_images {};
uint32_t sampler_variable = 0;
uint32_t sampler_type = 0;
uint32_t sampler_pointer_type = 0;
@@ -431,6 +432,8 @@ struct MemoryResourceAccess {
uint32_t EmitIndirectResourceIndex(EmitterState& state, uint32_t key, uint32_t mapping_offset,
uint32_t search_iterations, uint32_t default_resource);
// Slot word of an image heap byte offset; offsets past the heap select the null element.
uint32_t EmitImageHeapEntry(EmitterState& state, uint32_t key, uint32_t mapping_offset);
MemoryResourceAccess PrepareMemoryResourceAccess(EmitterState& state, const IR::MemoryInfo& mem,
uint32_t array_index = 0u);
@@ -38,19 +38,39 @@ uint32_t EmitBinaryU32(EmitterState& state, spv::Op opcode, uint32_t lhs, uint32
return ret;
}
static uint32_t LoadFlattenedSrt(EmitterState& state, uint32_t index) {
const auto pointer = state.builder.AllocateId();
state.builder.AddFunction(spv::OpAccessChain, TypeStorageBufferElementPointer(state), pointer,
state.flattened_srt_variable, ConstantU32(state, 0), index);
const auto value = state.builder.AllocateId();
state.builder.AddFunction(spv::OpLoad, TypeU32(state), value, pointer);
return value;
}
uint32_t EmitImageHeapEntry(EmitterState& state, uint32_t key, uint32_t mapping_offset) {
const auto Mapping = [&](uint32_t word) {
return LoadFlattenedSrt(state, ConstantU32(state, mapping_offset + word));
};
// Keys past the table select its trailing null slot.
const auto index = EmitUMin32(state, Binary(state, spv::OpShiftRightLogical, TypeU32(state), key,
ConstantU32(state, 2u)),
Mapping(0u));
const auto offset = Unary(state, spv::OpUConvert, TypeU64(state),
Binary(state, spv::OpShiftLeftLogical, TypeU32(state), index,
ConstantU32(state, 2u)));
const auto pointer = Unary(state, spv::OpConvertUToPtr, TypePhysicalU32Pointer(state),
Binary(state, spv::OpIAdd, TypeU64(state),
PackU64(state, Mapping(1u), Mapping(2u)), offset));
const auto entry = state.builder.AllocateId();
state.builder.AddFunction(spv::OpLoad, TypeU32(state), entry, pointer,
spv::MemoryAccessAlignedMask, 4u);
return entry;
}
uint32_t EmitIndirectResourceIndex(EmitterState& state, uint32_t key, uint32_t mapping_offset,
uint32_t search_iterations, uint32_t default_resource) {
const auto LoadMapping = [&](uint32_t index) {
const auto pointer = state.builder.AllocateId();
state.builder.AddFunction(spv::OpAccessChain, TypeStorageBufferElementPointer(state),
pointer, state.flattened_srt_variable, ConstantU32(state, 0),
index);
const auto value = state.builder.AllocateId();
state.builder.AddFunction(spv::OpLoad, TypeU32(state), value, pointer);
return value;
};
const auto mapping = ConstantU32(state, mapping_offset);
const auto count = LoadMapping(mapping);
const auto count = LoadFlattenedSrt(state, mapping);
if (search_iterations == 0u) {
const auto index = Binary(state, spv::OpShiftRightLogical, TypeU32(state), key,
ConstantU32(state, 2u));
@@ -58,7 +78,7 @@ uint32_t EmitIndirectResourceIndex(EmitterState& state, uint32_t key, uint32_t m
const auto safe = Select(state, TypeU32(state), valid, index, ConstantU32(state, 0u));
const auto entry = Binary(state, spv::OpIAdd, TypeU32(state), safe,
ConstantU32(state, mapping_offset + 1u));
return Select(state, TypeU32(state), valid, LoadMapping(entry),
return Select(state, TypeU32(state), valid, LoadFlattenedSrt(state, entry),
ConstantU32(state, default_resource));
}
const auto entry_at = [&](uint32_t index) {
@@ -76,7 +96,7 @@ uint32_t EmitIndirectResourceIndex(EmitterState& state, uint32_t key, uint32_t m
ConstantU32(state, 1u));
const auto probe = Select(state, TypeU32(state), active, mid, ConstantU32(state, 0u));
const auto less = Binary(state, spv::OpULessThan, TypeBool(state),
LoadMapping(entry_at(probe)), key);
LoadFlattenedSrt(state, entry_at(probe)), key);
low = Select(state, TypeU32(state),
Binary(state, spv::OpLogicalAnd, TypeBool(state), active, less),
Binary(state, spv::OpIAdd, TypeU32(state), mid, ConstantU32(state, 1u)), low);
@@ -87,9 +107,10 @@ uint32_t EmitIndirectResourceIndex(EmitterState& state, uint32_t key, uint32_t m
const auto in_range = Binary(state, spv::OpULessThan, TypeBool(state), low, count);
const auto entry = entry_at(Select(state, TypeU32(state), in_range, low, ConstantU32(state, 0u)));
const auto match = Binary(state, spv::OpLogicalAnd, TypeBool(state), in_range,
Binary(state, spv::OpIEqual, TypeBool(state), LoadMapping(entry), key));
const auto resource = LoadMapping(Binary(state, spv::OpIAdd, TypeU32(state), entry,
ConstantU32(state, 1u)));
Binary(state, spv::OpIEqual, TypeBool(state),
LoadFlattenedSrt(state, entry), key));
const auto resource = LoadFlattenedSrt(
state, Binary(state, spv::OpIAdd, TypeU32(state), entry, ConstantU32(state, 1u)));
return Select(state, TypeU32(state), match, resource, ConstantU32(state, default_resource));
}
@@ -187,7 +187,7 @@ void EmitMeshEntryPoint(EmitterState& state) {
vertex[component] = state.builder.AllocateId();
state.builder.AddFunction(
spv::OpBitFieldUExtract, TypeU32(state), vertex[component], packed,
ConstantU32(state, component * 10u), ConstantU32(state, 10));
ConstantU32(state, component * 10u), ConstantU32(state, 9));
}
const auto triangle = state.builder.AllocateId();
state.builder.AddFunction(spv::OpCompositeConstruct, TypeU32Vector(state, 3), triangle,
@@ -206,6 +206,17 @@ uint32_t F32ArrayType(EmitterState& state, uint32_t count) {
return state.builder.Type(spv::OpTypeArray, TypeF32(state), ConstantU32(state, count));
}
void DefineImageArray(EmitterState& state, ImageDefinition& definition, const IR::ImageResource& image,
uint32_t count, uint32_t set, uint32_t binding) {
definition.type = ImageType(state, image);
const auto array = state.builder.Type(spv::OpTypeArray, definition.type, ConstantU32(state, count));
definition.variable = state.builder.DefineGlobalVariable(
TypePointer(state, spv::StorageClassUniformConstant, array), spv::StorageClassUniformConstant);
state.builder.AddAnnotation(spv::OpDecorate, definition.variable, spv::DecorationDescriptorSet, set);
state.builder.AddAnnotation(spv::OpDecorate, definition.variable, spv::DecorationBinding, binding);
definition.pointer_type = TypePointer(state, spv::StorageClassUniformConstant, definition.type);
}
void DefineDescriptors(EmitterState& state) {
using Kind = IR::DescriptorBindingKind;
auto& info = state.program.info;
@@ -222,6 +233,14 @@ void DefineDescriptors(EmitterState& state) {
(image.mip_mode != IR::ImageMipMode::Dynamic && image.mip_count != 1u)) {
EXIT("shader image has an invalid descriptor type or mip count");
}
if (image.heap) {
// Heap kinds alias one array; the root itself is never accessed.
auto& definition = state.heap_images[IR::ImageBindingIndex(*kind)];
if (image.indirect_resources.empty() && definition.variable == 0) {
DefineImageArray(state, definition, image, IR::ImageHeapCapacity, IR::ImageHeapSet, 0u);
}
continue;
}
auto& count = counts[static_cast<size_t>(*kind)];
image.descriptor_index = count;
if (count == 0) image_types[IR::ImageBindingIndex(*kind)] = &image;
@@ -331,21 +350,11 @@ void DefineDescriptors(EmitterState& state) {
default: {
EXIT_IF(IR::ImageBindingResourceClass(kind) ==
IR::ImageResourceClass::None);
const auto& image = *image_types[IR::ImageBindingIndex(kind)];
const auto name = "image_" + std::to_string(static_cast<uint32_t>(kind));
auto& definition = state.images[IR::ImageBindingIndex(kind)];
definition.type = ImageType(state, image);
definition.variable = Define(ArrayType(definition.type), name.c_str(),
spv::StorageClassUniformConstant);
definition.pointer_type = TypePointer(
state, spv::StorageClassUniformConstant, definition.type);
if (image.dimension == ImageDimension::Dim1D ||
image.dimension == ImageDimension::Dim1DArray) {
state.builder.RequireCapability(image.resource_class ==
IR::ImageResourceClass::Sampled
? spv::CapabilitySampled1D
: spv::CapabilityImage1D);
}
DefineImageArray(state, definition, *image_types[IR::ImageBindingIndex(kind)],
counts[index], 0u, IR::NativeBinding(state.program.stage, kind));
state.builder.AddName(definition.variable,
("image_" + std::to_string(static_cast<uint32_t>(kind))).c_str());
break;
}
}
@@ -752,7 +761,7 @@ void DefineModule(EmitterState& state) {
if (state.program.info.buffer_u16) {
state.builder.RequireCapability(spv::CapabilityStorageBuffer16BitAccess);
}
if (state.program.info.uses_dma) {
if (state.program.info.uses_dma || state.program.info.uses_image_heap) {
state.builder.RequireCapability(spv::CapabilityInt64);
state.builder.RequireCapability(spv::CapabilityPhysicalStorageBufferAddresses);
state.builder.RequireExtension("SPV_KHR_physical_storage_buffer");
@@ -820,7 +829,7 @@ void DefineModule(EmitterState& state) {
state.builder.RequireExtension("SPV_KHR_fragment_shader_barycentric");
}
state.builder.RequireExtension("SPV_KHR_float_controls");
state.builder.AddMemoryModel(state.program.info.uses_dma
state.builder.AddMemoryModel(state.program.info.uses_dma || state.program.info.uses_image_heap
? spv::AddressingModelPhysicalStorageBuffer64
: spv::AddressingModelLogical,
spv::MemoryModelGLSL450);
@@ -235,9 +235,9 @@ void EmitDirectInstruction(ValueEmitContext& ctx, const IR::Inst& inst) {
const auto op = inst.GetOpcode();
if (ctx.half != 0 &&
(op == Op::Ballot || op == Op::ReadFirstLane || op == Op::SubgroupUMinU32 ||
op == Op::ReadConstBuffer ||
op == Op::ReadConstBuffer || op == Op::ReadRealtimeClock ||
(op == Op::LoadAddressU32 && ctx.Memory(inst).kind == IR::ResourceKind::ScalarAddress))) {
// Whole-wave results and SMEM loads (SGPR operands, no EXEC) are identical in both halves.
// Whole-wave results, clocks and SMEM loads (SGPR operands, no EXEC) match in both halves.
ctx.Define(inst, ctx.other_half->Result(inst));
return;
}
@@ -57,6 +57,8 @@ constexpr OpcodeMap SOP1_OPCODE_LIST[] = {
{0x14u, Opcode::S_FF1_I32_B64},
{0x15u, Opcode::S_FLBIT_I32_B32},
{0x16u, Opcode::S_FLBIT_I32_B64},
{0x17u, Opcode::S_FLBIT_I32},
{0x18u, Opcode::S_FLBIT_I32_I64},
{0x19u, Opcode::S_SEXT_I32_I8},
{0x1au, Opcode::S_SEXT_I32_I16},
{0x1bu, Opcode::S_BITSET0_B32},
@@ -60,6 +60,8 @@ enum class Opcode {
S_FF1_I32_B64,
S_FLBIT_I32_B32,
S_FLBIT_I32_B64,
S_FLBIT_I32,
S_FLBIT_I32_I64,
S_BITREPLICATE_B64_B32,
S_QUADMASK_B64,
S_GETPC_B64,
@@ -221,8 +221,17 @@ void Translator::V_XOR3_B32(const Decoder::Instruction& inst) {
WriteOperand(DestinationOperand(inst), result);
}
void Translator::S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb) {
void Translator::S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb,
bool sign) {
auto source = ReadU32Pair(inst.src0);
if (sign) {
// A signed source sign-extends a literal; then bits equal to the sign become zero.
if (inst.src0.kind == Decoder::OperandKind::LiteralConstant) {
source[1] = ir.ShiftRightArithmetic(source[0], IR::U32(IR::Value(31u)));
}
const auto mask = ir.ShiftRightArithmetic(source[1], IR::U32(IR::Value(31u)));
source = {ir.BitwiseXor(source[0], mask), ir.BitwiseXor(source[1], mask)};
}
for (auto& word: source) {
word = find_zero ? ir.BitwiseNot(word) : word;
}
@@ -240,12 +249,19 @@ void Translator::S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero,
void Translator::V_FFBH_32(const Decoder::Instruction& inst, bool sign) {
const auto source = ReadU32(inst.src0);
const auto value = sign ? ir.BitwiseXor(
source, ir.ShiftRightArithmetic(source, IR::U32(IR::Value(31u))))
: source;
const auto msb = IR::U32(ir.Emit(IR::ValueOpcode::FindUMsb32, {value}));
if (sign) {
// Bits equal to the sign become zero; FindILsb of the reversed bits is ~0 when none.
const auto bits =
ir.BitwiseXor(source, ir.ShiftRightArithmetic(source, IR::U32(IR::Value(31u))));
const auto reversed = ir.Emit(IR::ValueOpcode::BitReverse32, {bits});
WriteOperand(DestinationOperand(inst),
IR::U32(ir.Emit(IR::ValueOpcode::FindILsb32, {reversed})));
return;
}
// ResourceTracking matches this 31 - FindUMsb32 form.
const auto msb = IR::U32(ir.Emit(IR::ValueOpcode::FindUMsb32, {source}));
const auto position = ir.ISub(IR::U32(IR::Value(31u)), msb);
const auto result = ir.Select(ir.INotEqual(value, IR::U32(IR::Value(0u))), position,
const auto result = ir.Select(ir.INotEqual(source, IR::U32(IR::Value(0u))), position,
IR::U32(IR::Value(0xffffffffu)));
WriteOperand(DestinationOperand(inst), result);
}
@@ -1,11 +1,9 @@
#include "common/assert.h"
#include "common/logging/log.h"
#include "graphics/shader/recompiler/frontend/translate/Translator.h"
#include "graphics/shader/recompiler/frontend/decode/ImageOps.h"
#include <algorithm>
#include <array>
#include <atomic>
namespace Libs::Graphics::ShaderRecompiler::Frontend {
@@ -969,17 +967,9 @@ void Translator::EmitMemory(const Decoder::Instruction& inst) {
case Decoder::Opcode::S_BUFFER_LOAD_DWORDX4:
case Decoder::Opcode::S_BUFFER_LOAD_DWORDX8:
case Decoder::Opcode::S_BUFFER_LOAD_DWORDX16: return S_LOAD(inst, false);
case Decoder::Opcode::S_MEMREALTIME: {
static std::atomic_flag warned = ATOMIC_FLAG_INIT;
if (!warned.test_and_set(std::memory_order_relaxed)) {
Log::WriteToConsoleAndLog(
"Warning: S_MEMREALTIME uses placeholder UINT64_MAX; real-time clock not implemented.\n");
}
for (uint32_t component = 0; component < 2; component++) {
WriteOperand(ScalarDestinationOperand(inst.dst, component), IR::Value(UINT32_MAX));
}
return;
}
case Decoder::Opcode::S_MEMREALTIME:
return WriteU32Pair(inst.dst,
ExtractU64(IR::U64(ir.Emit(IR::ValueOpcode::ReadRealtimeClock))));
case Decoder::Opcode::BUFFER_LOAD_UBYTE:
case Decoder::Opcode::BUFFER_LOAD_UBYTE_D16:
@@ -238,10 +238,12 @@ void Translator::EmitScalar(const Decoder::Instruction& inst) {
return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseOr32, false, true, true);
case O::S_XNOR_B32:
return ComposedIntegerBinary(inst, IR::ValueOpcode::BitwiseXor32, false, true, true);
case O::S_FF0_I32_B64: return S_FF_I32_B64(inst, true, false);
case O::S_FF1_I32_B64: return S_FF_I32_B64(inst, false, false);
case O::S_FF0_I32_B64: return S_FF_I32_B64(inst, true, false, false);
case O::S_FF1_I32_B64: return S_FF_I32_B64(inst, false, false, false);
case O::S_FLBIT_I32_B32: return V_FFBH_32(inst, false);
case O::S_FLBIT_I32_B64: return S_FF_I32_B64(inst, false, true);
case O::S_FLBIT_I32_B64: return S_FF_I32_B64(inst, false, true, false);
case O::S_FLBIT_I32: return V_FFBH_32(inst, true);
case O::S_FLBIT_I32_I64: return S_FF_I32_B64(inst, false, true, true);
case O::S_BITSET0_B32: return S_BITSET_B32(inst, false);
case O::S_BITSET1_B32: return S_BITSET_B32(inst, true);
@@ -319,7 +319,10 @@ void Translator::WriteRawU32(const Decoder::Operand& operand, IR::U32 value) {
case Decoder::OperandKind::Vgpr: {
const auto reg = static_cast<IR::VectorReg>(operand.reg);
const auto old = ir.GetVectorReg(reg);
if (operand.dpp) {
// DPP skips EXEC lanes only through row/bank masks, or for invalid sources without
// BOUND_CTRL (dpp8 reads zero); otherwise it is a plain masked write.
if (operand.dpp && (operand.dpp_row_mask != 0xfu || operand.dpp_bank_mask != 0xfu ||
(!operand.dpp_bound_ctrl && !operand.dpp8))) {
value = IR::U32(ir.Emit(IR::ValueOpcode::DppUpdateU32, {value, old, ir.GetExec()},
DppFlags(operand)));
} else {
@@ -182,7 +182,7 @@ private:
void V_AND_OR_B32(const Decoder::Instruction& inst);
void V_OR3_B32(const Decoder::Instruction& inst);
void V_XOR3_B32(const Decoder::Instruction& inst);
void S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb);
void S_FF_I32_B64(const Decoder::Instruction& inst, bool find_zero, bool from_msb, bool sign);
void V_FFBH_32(const Decoder::Instruction& inst, bool sign);
void Integer24(const Decoder::Instruction& inst, bool sign, bool addend);
void V_MAD_U64_U32(const Decoder::Instruction& inst);
@@ -169,6 +169,7 @@ struct ImageResource {
bool depth_compare = false;
bool cube = false;
bool r128 = false;
bool heap = false; // Resident image heap root or kind.
uint32_t indirect_root = NoIndirectImage;
uint32_t indirect_mapping_offset = 0;
uint32_t indirect_search_iterations = 0;
@@ -344,6 +345,11 @@ inline constexpr uint32_t FirstImageBinding = 1u;
inline constexpr uint32_t FirstComparisonImageBinding = 22u;
inline constexpr uint32_t FirstStorageImageBinding = 29u;
inline constexpr uint32_t ImageBindingCount = 50u;
// Resident image heaps share one sampled-image array at binding 0 of this set. A heap slot
// word is kind << ImageHeapKindShift | element; element 0 is the null descriptor.
inline constexpr uint32_t ImageHeapSet = 1u;
inline constexpr uint32_t ImageHeapCapacity = 65536u;
inline constexpr uint32_t ImageHeapKindShift = 20u;
enum class DescriptorBindingKind : uint32_t {
Buffers = 0u,
@@ -529,6 +535,7 @@ struct ShaderInfo {
bool uses_lds = false;
bool uses_gds = false;
bool uses_flattened_srt = false;
bool uses_image_heap = false;
bool uses_dispatch_threads = false;
bool bvh = false;
bool subgroup_ballot = false;
@@ -10,6 +10,7 @@ VALUE_OPCODE(ConditionRef, U1, U1)
VALUE_OPCODE(GetUserData, U32, ScalarReg)
VALUE_OPCODE(GetDispatchThreadExtent, U32, U32)
VALUE_OPCODE(GetShaderBase, U64)
VALUE_OPCODE(ReadRealtimeClock, U64)
VALUE_OPCODE(MeshDrawParameter, U32, U32)
VALUE_OPCODE(MeshAllocate, Void, U32)
VALUE_OPCODE(TessellationBase, U32, U32)
@@ -681,6 +681,32 @@ bool MaterializeIndirectDescriptor(const ResourcePlan& p
const auto count = table_size >= first + 4u
? std::min((table_size - first - 4u) / step + 1u, domain)
: 0u;
descriptors[resource_index] = {.dword_count = dword_count};
if constexpr (std::is_same_v<Specialization, ResourceSpecialization::Image>) {
const auto& image = program.info.images[resource_index];
// A resident heap maps slot words instead of materializing its entries.
ImageHeapView view;
if (image_heap && runtime.image_heap != nullptr &&
image.resource_class == ImageResourceClass::Sampled &&
image.mip_mode != ImageMipMode::Dynamic &&
runtime.image_heap->Acquire({table_base, table_size, step, first, count, &image},
view)) {
snapshot.specialization_reads.emplace_back(table_base, table_size);
if (view.kinds.empty()) return true;
auto root = root_resource;
root.indirect_root = resource_index;
root.indirect_mapping_offset = static_cast<uint32_t>(snapshot.flattened_srt.size());
root.heap = true;
specializations[resource_index] = root;
snapshot.flattened_srt.insert(snapshot.flattened_srt.end(),
{view.slot_count, static_cast<uint32_t>(view.slots), static_cast<uint32_t>(view.slots >> 32u)});
for (const auto& kind: view.kinds) {
descriptors.push_back(kind);
specializations.push_back(root);
}
return true;
}
}
if (child_count == 0u || (!image_heap && count * child_count > MaxIndirectDescriptorProbes))
return failed("candidate count", fmt::format("rows={} limit={}", count, MaxIndirectDescriptorProbes));
// Affine SMEM offsets wrap in U32 before adding the aligned immediate.
@@ -692,7 +718,6 @@ bool MaterializeIndirectDescriptor(const ResourcePlan& p
return read_failed("image heap read", 0u, 0u, table_base, 0u, heap_words.size() * sizeof(uint32_t));
}
const auto mapping_offset = snapshot.flattened_srt.size();
descriptors[resource_index] = {.dword_count = dword_count};
snapshot.flattened_srt.resize(mapping_offset + 2u, 0u);
for (uint32_t word = 0; word < count; ++word) {
const auto table_offset = affine ? first + word * step : uint64_t {word} * 4u;
@@ -1010,10 +1035,12 @@ static bool BuildResourceSpecialization(const ResourcePlan& program, ResourceSna
const auto key_count = root.indirect_mapping_offset < snapshot.flattened_srt.size()
? snapshot.flattened_srt[root.indirect_mapping_offset]
: 0u;
// A heap maps its slot count and slot address; tables map each key.
size_t mapping_words = 3u;
if (!root.heap)
mapping_words = 1u + size_t {key_count} * (root.indirect_search_iterations == 0u ? 1u : 2u);
if (key_count == 0u ||
static_cast<size_t>(root.indirect_mapping_offset) + 1u +
static_cast<size_t>(key_count) * (root.indirect_search_iterations == 0u ? 1u : 2u) >
snapshot.flattened_srt.size()) {
root.indirect_mapping_offset + mapping_words > snapshot.flattened_srt.size()) {
return SpecializationFail("indirect image specialization has an invalid key mapping");
}
uint32_t exemplar = ImageResource::NoIndirectImage;
@@ -1749,6 +1776,43 @@ bool MaterializeResources(const ResourcePlan& program, const SrtRuntime& runtime
return BuildResourceSpecialization(program, snapshot, specialization);
}
uint32_t ImageHeapKind(const ImageResource& root, DescriptorValue& value) {
// Mirrors BuildResourceSpecialization for sampled roots: entries it would reject, and
// entries that need shader-side conversion, read as null instead of failing the draw.
const auto format = static_cast<Prospero::BufferFormat>((value.dwords[1] >> 20u) & 0x1ffu);
const auto tile = static_cast<Prospero::TileMode>((value.dwords[3] >> 20u) & 0x1fu);
const auto numeric_class =
Prospero::SampledTextureNumericClass(Prospero::RemapTextureFormat(format));
if (NullImageDescriptor(value) || !ValidImageDescriptor(value, root.r128) ||
DescriptorDimension(value, root.dimension) != root.dimension ||
Prospero::IsFmaskTextureFormat(format) || format == Prospero::BufferFormat::k11_11_10UNorm ||
ImageConversionFormat(format) != Prospero::BufferFormat::kInvalid ||
(root.unfiltered_access_only && format == Prospero::BufferFormat::k16UInt &&
tile == Prospero::TileMode::kDepth) ||
numeric_class == Prospero::TextureNumericClass::Unsupported ||
(root.depth_compare && numeric_class != Prospero::TextureNumericClass::Float)) {
value.dwords.fill(0u);
return 0u;
}
return 1u + static_cast<uint32_t>(numeric_class) * 2u + (DescriptorIsCube(value) ? 1u : 0u);
}
ImageResource ImageHeapResource(const ImageResource& root, uint32_t kind) {
ImageResource resource;
resource.resource_class = ImageResourceClass::Sampled;
resource.read = true;
resource.heap = true;
resource.dimension = root.dimension;
resource.unfiltered_access_only = root.unfiltered_access_only;
resource.depth_compare = root.depth_compare;
resource.r128 = root.r128;
if (kind != 0u) {
resource.numeric_class = static_cast<Prospero::TextureNumericClass>((kind - 1u) / 2u);
resource.cube = ((kind - 1u) & 1u) != 0u;
}
return resource;
}
void ApplyResourceSpecialization(Program& program, const ResourceSpecialization& specialization) {
EXIT_IF(!program.resource_tracking_complete || program.shader_info_complete);
EXIT_IF(program.info.buffers.size() > specialization.buffers.size() ||
@@ -1799,6 +1863,7 @@ void ApplyResourceSpecialization(Program& program, const ResourceSpecialization&
image.indirect_mapping_offset = source.indirect_mapping_offset;
image.indirect_search_iterations = source.indirect_search_iterations;
image.cube = source.cube;
image.heap = source.heap;
image.indirect_resources.clear();
}
for (uint32_t index = 0; index < images.size(); index++) {
@@ -31,6 +31,7 @@ struct ResourceSpecialization {
uint32_t indirect_search_iterations = 0;
bool cube = false;
bool fmask = false;
bool heap = false;
bool operator==(const Image&) const = default;
};
@@ -59,6 +60,13 @@ bool MaterializeResources(const ResourcePlan& program, const SrtRuntime& runtime
// Applies an already-derived specialization to native IR before layout and emission.
void ApplyResourceSpecialization(Program& program, const ResourceSpecialization& specialization);
// Kind of an image heap entry as root sees it, or 0 (and a zeroed value) when root reads it as
// null. Entries of one kind specialize identically.
uint32_t ImageHeapKind(const ImageResource& root, DescriptorValue& value);
// The sampled resource that root resolves entries of kind as; kind 0 gives the identity of
// root's view of a heap.
ImageResource ImageHeapResource(const ImageResource& root, uint32_t kind);
} // namespace Libs::Graphics::ShaderRecompiler::IR
#endif /* EMULATOR_INCLUDE_EMULATOR_GRAPHICS_SHADER_RECOMPILER_RESOURCEMATERIALIZATION_H_ */
@@ -594,6 +594,7 @@ void CollectShaderInfo(Program& program, ShaderStageInputInfo input_info) {
};
info.uses_flattened_srt |= std::ranges::any_of(info.buffers, uses_mapping) ||
std::ranges::any_of(info.images, uses_mapping);
info.uses_image_heap = std::ranges::any_of(info.images, &ImageResource::heap);
switch (program.stage) {
case ShaderType::Vertex:
case ShaderType::Local: CollectVertexInputs(input_info.vertex, info, inputs.components); break;
@@ -12,6 +12,33 @@ class Value;
using SrtMemoryReader = bool (*)(void* userdata, uint64_t address, std::span<uint32_t> values);
// Image descriptors at base + first + k * step, k < count, within size bytes that a sampled
// root selects by GPU-computed byte offset. The host keeps them resident instead of
// materializing them.
struct ImageHeapRequest {
uint64_t base = 0;
uint64_t size = 0;
uint64_t step = 0;
uint64_t first = 0;
uint64_t count = 0;
const ImageResource* root = nullptr;
};
struct ImageHeapView {
uint64_t slots = 0; // Device address of one slot per heap dword.
uint32_t slot_count = 0; // Keys past it select a trailing null slot.
std::span<const DescriptorValue> kinds; // One representative descriptor per kind.
};
class ImageHeapSource {
public:
// False when the host does not keep the table resident.
virtual bool Acquire(const ImageHeapRequest& request, ImageHeapView& view) = 0;
protected:
~ImageHeapSource() = default;
};
struct SrtRuntime {
std::span<const uint32_t> user_data;
uint64_t shader_base = 0;
@@ -20,6 +47,7 @@ struct SrtRuntime {
SrtMemoryReader read_specialization_memory = nullptr;
std::span<const uint32_t> workgroup_counts;
uint64_t indirect_args_address = 0;
ImageHeapSource* image_heap = nullptr;
};
class SrtReadCapture {
+37 -98
View File
@@ -688,7 +688,8 @@ public:
bool Find(uint64_t phys_addr, bool next, PhysicalMemory::AllocatedBlock* out);
bool CanMapDirect(uint64_t phys_addr, size_t len);
bool ReleasePoolExpansion(uint64_t phys_addr, size_t len);
bool GetAllocatedSpan(uint64_t phys_addr, size_t len, std::vector<AllocatedBlock>* blocks);
bool GetAllocatedSpan(uint64_t phys_addr, size_t len, bool skip_gaps,
std::vector<AllocatedBlock>* blocks);
std::vector<AllocatedBlock> FindMappings(uint64_t phys_addr, size_t len);
void ProtectMapping(uint64_t vaddr, uint64_t size, int prot, VirtualMemory::Mode mode,
GpuAccessMode gpu_mode);
@@ -1353,87 +1354,31 @@ bool PhysicalMemory::Map(uint64_t vaddr, uint64_t phys_addr, size_t len, int pro
int memory_type) {
Common::LockGuard lock(m_mutex);
if (len == 0 || UINT64_MAX - phys_addr < len) {
if (!CanMapDirect(phys_addr, len)) {
return false;
}
auto current = phys_addr;
auto next = m_physical.upper_bound(current);
if (next == m_physical.begin()) {
return false;
}
auto first = std::prev(next);
bool has_automatic = false;
while (current < phys_addr + len) {
auto block = m_physical.upper_bound(current);
if (block == m_physical.begin()) {
return false;
}
--block;
const auto block_end = block->second.start_addr + block->second.size;
if (block->second.kind == AllocationKind::Pooled || current < block->second.start_addr ||
current >= block_end) {
return false;
}
has_automatic |= block->second.kind == AllocationKind::Automatic;
current = std::min<uint64_t>(phys_addr + len, block_end);
if (current < phys_addr + len) {
const auto following = std::next(block);
if (following == m_physical.end() || following->second.start_addr != current) {
return false;
}
}
}
AllocatedBlock mapping = first->second;
mapping.start_addr = phys_addr;
mapping.size = len;
mapping.map_vaddr = vaddr;
mapping.map_size = len;
mapping.host_vaddr = vaddr;
mapping.host_size = len;
mapping.prot = prot;
mapping.mode = mode;
mapping.gpu_mode = gpu_mode;
mapping.memory_type = memory_type;
AllocatedBlock mapping {};
mapping.start_addr = phys_addr;
mapping.size = len;
mapping.map_vaddr = vaddr;
mapping.map_size = len;
mapping.host_vaddr = vaddr;
mapping.host_size = len;
mapping.prot = prot;
mapping.mode = mode;
mapping.gpu_mode = gpu_mode;
mapping.memory_type = memory_type;
CopyVirtualRangeName(mapping.name, name);
m_mappings.push_back(mapping);
if (has_automatic) {
RemoveFreeRange(m_automatic_free, phys_addr, len);
}
RemoveFreeRange(m_automatic_free, phys_addr, len);
return true;
}
bool PhysicalMemory::CanMapDirect(uint64_t phys_addr, size_t len) {
Common::LockGuard lock(m_mutex);
if (len == 0 || UINT64_MAX - phys_addr < len) {
return false;
}
const auto end = phys_addr + len;
auto current = phys_addr;
while (current < end) {
auto block = m_physical.upper_bound(current);
if (block == m_physical.begin()) {
return false;
}
--block;
const auto block_end = block->second.start_addr + block->second.size;
if (block->second.kind == AllocationKind::Pooled || current < block->second.start_addr ||
current >= block_end) {
return false;
}
current = std::min(end, block_end);
if (current < end) {
const auto following = std::next(block);
if (following == m_physical.end() || following->second.start_addr != current) {
return false;
}
}
}
return true;
std::vector<AllocatedBlock> blocks;
return GetAllocatedSpan(phys_addr, len, false, &blocks);
}
bool PhysicalMemory::ReleasePoolExpansion(uint64_t phys_addr, size_t len) {
@@ -1448,7 +1393,7 @@ bool PhysicalMemory::ReleasePoolExpansion(uint64_t phys_addr, size_t len) {
return true;
}
bool PhysicalMemory::GetAllocatedSpan(uint64_t phys_addr, size_t len,
bool PhysicalMemory::GetAllocatedSpan(uint64_t phys_addr, size_t len, bool skip_gaps,
std::vector<AllocatedBlock>* blocks) {
EXIT_IF(blocks == nullptr);
blocks->clear();
@@ -1459,34 +1404,26 @@ bool PhysicalMemory::GetAllocatedSpan(uint64_t phys_addr, size_t len,
Common::LockGuard lock(m_mutex);
const auto end = phys_addr + len;
auto current = phys_addr;
while (current < end) {
auto block = m_physical.upper_bound(current);
if (block == m_physical.begin()) {
blocks->clear();
return false;
}
auto block = m_physical.upper_bound(phys_addr);
if (block != m_physical.begin()) {
--block;
}
for (; block != m_physical.end() && block->second.start_addr < end; ++block) {
const auto block_end = block->second.start_addr + block->second.size;
if (block->second.kind == AllocationKind::Pooled || current < block->second.start_addr ||
current >= block_end) {
blocks->clear();
if (block_end <= current) {
continue;
}
if (block->second.kind == AllocationKind::Pooled ||
(!skip_gaps && current < block->second.start_addr)) {
return false;
}
auto part = block->second;
part.start_addr = current;
part.size = std::min(end, block_end) - current;
part.start_addr = std::max(current, part.start_addr);
current = std::min(end, block_end);
part.size = current - part.start_addr;
blocks->push_back(part);
current += part.size;
if (current < end) {
const auto following = std::next(block);
if (following == m_physical.end() || following->second.start_addr != current) {
blocks->clear();
return false;
}
}
}
return true;
return skip_gaps || current == end;
}
bool PhysicalMemory::Unmap(uint64_t vaddr, uint64_t size, GpuAccessMode* gpu_mode,
@@ -2855,7 +2792,7 @@ int KYTY_SYSV_ABI KernelAllocateMainDirectMemory(size_t len, size_t alignment, i
alignment, memory_type, phys_addr_out);
}
static int ReleaseDirectMemoryInternal(int64_t start, size_t len) {
static int ReleaseDirectMemoryInternal(int64_t start, size_t len, bool skip_gaps) {
std::lock_guard<std::recursive_mutex> memory_operation_lock(g_memory_operation_mutex);
if (g_pooled_memory->ReleaseExpansion(static_cast<uint64_t>(start), len)) {
@@ -2866,7 +2803,8 @@ static int ReleaseDirectMemoryInternal(int64_t start, size_t len) {
}
std::vector<PhysicalMemory::AllocatedBlock> allocated_span;
if (!g_physical_memory->GetAllocatedSpan(static_cast<uint64_t>(start), len, &allocated_span)) {
if (!g_physical_memory->GetAllocatedSpan(static_cast<uint64_t>(start), len, skip_gaps,
&allocated_span)) {
return KERNEL_ERROR_EACCES;
}
@@ -2965,7 +2903,8 @@ int KYTY_SYSV_ABI KernelReleaseDirectMemory(int64_t start, size_t len) {
return validation;
}
if (len != 0) {
(void)ReleaseDirectMemoryInternal(start, len);
// An unchecked release frees the allocated parts and skips the unallocated ones.
(void)ReleaseDirectMemoryInternal(start, len, true);
}
return OK;
}
@@ -2982,7 +2921,7 @@ int KYTY_SYSV_ABI KernelCheckedReleaseDirectMemory(int64_t start, size_t len) {
return validation;
}
const int result = ReleaseDirectMemoryInternal(start, len);
const int result = ReleaseDirectMemoryInternal(start, len, false);
return result == KERNEL_ERROR_EACCES ? KERNEL_ERROR_ENOENT : result;
}
+10
View File
@@ -37,6 +37,16 @@ endif()
target_include_directories(launcher PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/include" "${CMAKE_CURRENT_SOURCE_DIR}/src")
set_target_properties(launcher PROPERTIES AUTOUIC ON AUTOMOC ON AUTORCC ON)
add_executable(launcher_game_views_tests EXCLUDE_FROM_ALL
"${KYTY_TESTS_DIR}/LauncherGameViewsTests.cpp"
src/gameContent.cpp
src/launcherTheme.cpp
resources.qrc
)
target_link_libraries(launcher_game_views_tests Qt6::Widgets common)
target_include_directories(launcher_game_views_tests PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/include")
set_target_properties(launcher_game_views_tests PROPERTIES AUTOUIC ON AUTORCC ON)
if(WIN32 AND NOT KYTY_CLANG_CL)
set_target_properties(launcher PROPERTIES LINK_FLAGS "${KYTY_LD_OPTIONS}")
elseif(LINUX)
@@ -180,6 +180,12 @@
</widget>
</item>
<item row="8" column="1">
<widget class="QCheckBox" name="checkBox_occlusion_queries">
<property name="toolTip"><string>Count the samples of occlusion queries on the GPU, so hidden lights and objects stay hidden. Costs some performance.</string></property>
<property name="text"><string>Enable occlusion queries</string></property>
</widget>
</item>
<item row="9" column="1">
<widget class="QCheckBox" name="checkBox_hdr">
<property name="toolTip"><string>Enable HDR10 output for supported games. Requires HDR enabled on the host display.</string></property>
<property name="text"><string>Enable HDR display support</string></property>
@@ -304,8 +304,11 @@
</layout>
</item>
<item>
<widget class="QStackedWidget" name="game_views">
<widget class="GameListTreeWidget" name="cfgs_list">
<widget class="GameViews" name="game_views">
<widget class="QTreeWidget" name="cfgs_list">
<property name="styleSheet">
<string notr="true">#cfgs_list { background: transparent; }</string>
</property>
<property name="rootIsDecorated">
<bool>false</bool>
</property>
@@ -350,7 +353,11 @@
</property>
</column>
</widget>
<widget class="QListView" name="cfgs_grid"/>
<widget class="QListView" name="cfgs_grid">
<property name="styleSheet">
<string notr="true">#cfgs_grid { background: transparent; }</string>
</property>
</widget>
</widget>
</item>
</layout>
@@ -369,9 +376,10 @@
</widget>
<customwidgets>
<customwidget>
<class>GameListTreeWidget</class>
<extends>QTreeWidget</extends>
<header>gameListTreeWidget.h</header>
<class>GameViews</class>
<extends>QStackedWidget</extends>
<header>gameViews.h</header>
<container>1</container>
</customwidget>
</customwidgets>
<resources/>
+2
View File
@@ -129,6 +129,7 @@ public:
bool readback_tiled_images = false;
bool sync_raw_image_buffers = false;
bool tessellation_enabled = false;
bool occlusion_queries_enabled = false;
bool trophy_enabled = true;
int trophy_duration_seconds = 6;
bool perf_overlay_enabled = false;
@@ -169,6 +170,7 @@ public:
readback_tiled_images = other.readback_tiled_images;
sync_raw_image_buffers = other.sync_raw_image_buffers;
tessellation_enabled = other.tessellation_enabled;
occlusion_queries_enabled = other.occlusion_queries_enabled;
trophy_enabled = other.trophy_enabled;
trophy_duration_seconds = other.trophy_duration_seconds;
perf_overlay_enabled = other.perf_overlay_enabled;
-77
View File
@@ -1,77 +0,0 @@
#ifndef GAME_LIST_TREE_WIDGET_H
#define GAME_LIST_TREE_WIDGET_H
#include "gameContent.h"
#include <QPaintEvent>
#include <QPainter>
#include <QPalette>
#include <QPixmap>
#include <QPoint>
#include <QResizeEvent>
#include <QString>
#include <QTreeWidget>
class GameListTreeWidget: public QTreeWidget {
public:
explicit GameListTreeWidget(QWidget* parent = nullptr): QTreeWidget(parent) {
setAutoFillBackground(false);
viewport()->setAutoFillBackground(false);
}
void SetBackgroundImage(const QString& key) {
if (m_path == key) {
return;
}
m_path = key;
m_source = QPixmap();
if (!key.isEmpty()) {
m_source.loadFromData(GameContent::ReadFile(key, QStringLiteral("sce_sys/pic0.png"),
GameContent::MaxImageSize));
}
UpdateScaledBackground();
viewport()->update();
}
protected:
void resizeEvent(QResizeEvent* event) override {
QTreeWidget::resizeEvent(event);
UpdateScaledBackground();
}
void paintEvent(QPaintEvent* event) override {
{
QPainter painter(viewport());
if (m_scaled.isNull()) {
painter.fillRect(event->rect(), palette().brush(QPalette::Base));
} else {
painter.drawPixmap(m_scaled_pos, m_scaled);
auto overlay = palette().color(QPalette::Base);
overlay.setAlpha(185);
painter.fillRect(event->rect(), overlay);
}
}
QTreeWidget::paintEvent(event);
}
private:
void UpdateScaledBackground() {
m_scaled = QPixmap();
if (m_source.isNull() || viewport()->size().isEmpty()) {
return;
}
m_scaled = m_source.scaled(viewport()->size(), Qt::KeepAspectRatioByExpanding,
Qt::FastTransformation);
m_scaled_pos = QPoint((viewport()->width() - m_scaled.width()) / 2,
(viewport()->height() - m_scaled.height()) / 2);
}
QString m_path;
QPixmap m_source;
QPixmap m_scaled;
QPoint m_scaled_pos;
};
#endif // GAME_LIST_TREE_WIDGET_H
+148
View File
@@ -0,0 +1,148 @@
#ifndef GAME_VIEWS_H
#define GAME_VIEWS_H
#include "gameContent.h"
#include <QGuiApplication>
#include <QIcon>
#include <QListView>
#include <QPaintEvent>
#include <QPainter>
#include <QPalette>
#include <QPixmap>
#include <QResizeEvent>
#include <QStackedWidget>
#include <QString>
#include <QStyledItemDelegate>
#include <QTreeWidget>
inline constexpr int GAME_NAME_COLUMN = 0;
inline constexpr int GAME_SERIAL_COLUMN = 1;
inline constexpr int GAME_VERSION_COLUMN = 2;
inline constexpr int GAME_FIRMWARE_VERSION_COLUMN = 3;
inline constexpr int GAME_SIZE_COLUMN = 4;
inline constexpr int GAME_PATH_COLUMN = 5;
inline constexpr int GAME_STATUS_COLUMN = 6;
inline constexpr int GAME_COMMENT_COLUMN = 7;
// Holds the transparent list and grid views and paints the selected game's pic0 behind both.
class GameViews: public QStackedWidget {
public:
using QStackedWidget::QStackedWidget;
void SetBackgroundImage(const QString& key) {
if (m_path == key) {
return;
}
m_path = key;
m_source = QPixmap();
if (!key.isEmpty()) {
m_source.loadFromData(GameContent::ReadFile(key, QStringLiteral("sce_sys/pic0.png"),
GameContent::MaxImageSize));
}
UpdateScaledBackground();
update();
}
protected:
void resizeEvent(QResizeEvent* event) override {
QStackedWidget::resizeEvent(event);
UpdateScaledBackground();
}
void paintEvent(QPaintEvent* event) override {
QPainter painter(this);
auto overlay = palette().color(QPalette::Base);
if (!m_scaled.isNull()) {
painter.drawPixmap((width() - m_scaled.width()) / 2, (height() - m_scaled.height()) / 2,
m_scaled);
overlay.setAlpha(160);
}
painter.fillRect(event->rect(), overlay);
}
private:
void UpdateScaledBackground() {
m_scaled = m_source.isNull() ? QPixmap()
: m_source.scaled(size(), Qt::KeepAspectRatioByExpanding,
Qt::FastTransformation);
}
QString m_path;
QPixmap m_source;
QPixmap m_scaled;
};
// A card on a translucent panel, so the text keeps its contrast over the game background.
class GameGridDelegate: public QStyledItemDelegate {
public:
using QStyledItemDelegate::QStyledItemDelegate;
QSize sizeHint(const QStyleOptionViewItem& option,
const QModelIndex& /*index*/) const override {
return option.decorationSize +
QSize(2 * CARD_PADDING, 3 * CARD_PADDING + 2 * option.fontMetrics.lineSpacing());
}
void paint(QPainter* painter, const QStyleOptionViewItem& option,
const QModelIndex& index) const override {
// The stylesheet made the view's own palette transparent, so take the panel from the app's.
auto panel = QGuiApplication::palette().color(QPalette::Window);
auto text_color = option.palette.color(QPalette::Text);
panel.setAlpha(210);
painter->save();
painter->setRenderHint(QPainter::Antialiasing);
painter->setPen(option.state.testFlag(QStyle::State_Selected) ? QPen(text_color)
: Qt::NoPen);
painter->setBrush(panel);
painter->drawRoundedRect(QRectF(option.rect).adjusted(0.5, 0.5, -0.5, -0.5), 4, 4);
const QRect cover(option.rect.topLeft() + QPoint(CARD_PADDING, CARD_PADDING),
option.decorationSize);
index.data(Qt::DecorationRole).value<QIcon>().paint(painter, cover);
QRect line(cover.left(), cover.bottom() + 1 + CARD_PADDING, cover.width(),
option.fontMetrics.lineSpacing());
auto title_font = option.font;
title_font.setBold(true);
painter->setFont(title_font);
painter->setPen(text_color);
painter->drawText(line, Qt::AlignLeft | Qt::AlignVCenter,
painter->fontMetrics().elidedText(index.data().toString(), Qt::ElideRight,
line.width()));
line.translate(0, line.height());
const auto size = index.siblingAtColumn(GAME_SIZE_COLUMN).data().toString();
const auto details = QStringLiteral("%1 %2").arg(
index.siblingAtColumn(GAME_SERIAL_COLUMN).data().toString(),
index.siblingAtColumn(GAME_VERSION_COLUMN).data().toString());
text_color.setAlpha(200);
painter->setFont(option.font);
painter->setPen(text_color);
painter->drawText(line, Qt::AlignRight | Qt::AlignVCenter, size);
line.setRight(line.right() - option.fontMetrics.horizontalAdvance(size) - CARD_PADDING);
painter->drawText(line, Qt::AlignLeft | Qt::AlignVCenter,
option.fontMetrics.elidedText(details, Qt::ElideRight, line.width()));
painter->restore();
}
private:
static constexpr int CARD_PADDING = 8;
};
// Shows the list's games as cards in the grid: both views share the items and the current game.
inline void SetupGameGrid(QTreeWidget* list, QListView* grid) {
grid->setModel(list->model());
grid->setSelectionModel(list->selectionModel());
grid->setViewMode(QListView::IconMode);
grid->setMovement(QListView::Static);
grid->setResizeMode(QListView::Adjust);
grid->setSpacing(8);
grid->setUniformItemSizes(true);
grid->setItemDelegate(new GameGridDelegate(grid));
// A card also shows other columns (the size arrives late), which the grid does not repaint on.
QObject::connect(list->model(), &QAbstractItemModel::dataChanged, grid->viewport(),
qOverload<>(&QWidget::update));
}
#endif // GAME_VIEWS_H
+2
View File
@@ -16,6 +16,7 @@ QVariantMap Configuration::GameSettings() const {
KYTY_CFG_SET(readback_tiled_images);
KYTY_CFG_SET(sync_raw_image_buffers);
KYTY_CFG_SET(tessellation_enabled);
KYTY_CFG_SET(occlusion_queries_enabled);
KYTY_CFG_SET(trophy_enabled);
KYTY_CFG_SET(trophy_duration_seconds);
KYTY_CFG_SET(perf_overlay_enabled);
@@ -66,6 +67,7 @@ void Configuration::ReadGameSettingsValues(const Settings& s) {
KYTY_CFG_GET(readback_tiled_images);
KYTY_CFG_GET(sync_raw_image_buffers);
KYTY_CFG_GET(tessellation_enabled);
KYTY_CFG_GET(occlusion_queries_enabled);
trophy_enabled = s.value("trophy_enabled", trophy_enabled).toBool();
KYTY_CFG_GET(trophy_duration_seconds);
KYTY_CFG_GET(perf_overlay_enabled);
@@ -289,6 +289,7 @@ void ConfigurationEditDialog::Init(const Configuration& info) {
m_ui->checkBox_readback_tiled->setChecked(info.readback_tiled_images);
m_ui->checkBox_sync_raw_image_buffers->setChecked(info.sync_raw_image_buffers);
m_ui->checkBox_tessellation->setChecked(info.tessellation_enabled);
m_ui->checkBox_occlusion_queries->setChecked(info.occlusion_queries_enabled);
m_ui->checkBox_skip_notice_screen->setChecked(info.skip_notice_screen);
m_ui->checkBox_network_connected->setChecked(info.network_connected);
m_ui->checkBox_trophy_notifications->setChecked(info.trophy_enabled);
@@ -483,6 +484,7 @@ static void UpdateInfo(Configuration& info, Ui::ConfigurationEditDialog& ui, boo
info.readback_tiled_images = ui.checkBox_readback_tiled->isChecked();
info.sync_raw_image_buffers = ui.checkBox_sync_raw_image_buffers->isChecked();
info.tessellation_enabled = ui.checkBox_tessellation->isChecked();
info.occlusion_queries_enabled = ui.checkBox_occlusion_queries->isChecked();
info.skip_notice_screen = ui.checkBox_skip_notice_screen->isChecked();
info.network_connected = ui.checkBox_network_connected->isChecked();
info.trophy_enabled = ui.checkBox_trophy_notifications->isChecked();
+11 -71
View File
@@ -8,7 +8,7 @@
#include "configurationEditDialog.h"
#include "configurationItem.h"
#include "gameContent.h"
#include "gameListTreeWidget.h"
#include "gameViews.h"
#include "inputMappingDialog.h"
#include "mainDialog.h"
#include "patchesDialog.h"
@@ -53,7 +53,6 @@
#include <QSignalBlocker>
#include <QSlider>
#include <QStyle>
#include <QStyledItemDelegate>
#include <QToolButton>
#include <QTreeWidget>
#include <QUrl>
@@ -83,15 +82,6 @@ constexpr char CONF_ADDONS_DIR[] = "addons_dir";
constexpr char CONF_GLOBAL[] = "GlobalConfiguration";
constexpr char SAVE_DATA_DIR[] = "_SaveData";
constexpr int GAME_NAME_COLUMN = 0;
constexpr int GAME_SERIAL_COLUMN = 1;
constexpr int GAME_VERSION_COLUMN = 2;
constexpr int GAME_FIRMWARE_VERSION_COLUMN = 3;
constexpr int GAME_SIZE_COLUMN = 4;
constexpr int GAME_PATH_COLUMN = 5;
constexpr int GAME_STATUS_COLUMN = 6;
constexpr int GAME_COMMENT_COLUMN = 7;
static QString NormalizeGameDirectory(const QString& dir) {
const auto trimmed = dir.trimmed();
if (trimmed.isEmpty()) {
@@ -151,54 +141,6 @@ static QStringList SettingsStringList(const QVariant& value) {
return text.isEmpty() ? QStringList() : QStringList({text});
}
class GameGridDelegate: public QStyledItemDelegate {
public:
using QStyledItemDelegate::QStyledItemDelegate;
QSize sizeHint(const QStyleOptionViewItem& option, const QModelIndex&) const override {
return option.decorationSize + QSize(16, 24 + option.fontMetrics.lineSpacing() * 2);
}
protected:
void initStyleOption(QStyleOptionViewItem* option, const QModelIndex& index) const override {
const auto image_size = option->decorationSize;
QStyledItemDelegate::initStyleOption(option, index);
option->decorationSize = image_size;
option->decorationPosition = QStyleOptionViewItem::Top;
option->decorationAlignment = Qt::AlignCenter;
option->displayAlignment = Qt::AlignHCenter | Qt::AlignTop;
option->features |= QStyleOptionViewItem::WrapText;
}
};
static void ConfigureGameList(Ui::ConfigurationListWidget* ui) {
ui->cfgs_list->setAlternatingRowColors(false);
ui->cfgs_list->setAllColumnsShowFocus(true);
ui->cfgs_list->setIndentation(0);
ui->cfgs_list->setMouseTracking(false);
ui->cfgs_list->setSelectionBehavior(QAbstractItemView::SelectRows);
ui->cfgs_list->setSelectionMode(QAbstractItemView::SingleSelection);
ui->cfgs_list->setTextElideMode(Qt::ElideMiddle);
ui->cfgs_list->header()->setDefaultAlignment(Qt::AlignLeft | Qt::AlignVCenter);
ui->cfgs_list->header()->setHighlightSections(false);
ui->cfgs_list->header()->setStretchLastSection(true);
// Both presentations use the same items, icons and current game.
ui->cfgs_grid->setModel(ui->cfgs_list->model());
ui->cfgs_grid->setSelectionModel(ui->cfgs_list->selectionModel());
ui->cfgs_grid->setViewMode(QListView::IconMode);
ui->cfgs_grid->setMovement(QListView::Static);
ui->cfgs_grid->setResizeMode(QListView::Adjust);
ui->cfgs_grid->setSelectionMode(QAbstractItemView::SingleSelection);
ui->cfgs_grid->setSelectionBehavior(QAbstractItemView::SelectRows);
ui->cfgs_grid->setEditTriggers(QAbstractItemView::NoEditTriggers);
ui->cfgs_grid->setSpacing(8);
ui->cfgs_grid->setUniformItemSizes(true);
ui->cfgs_grid->setWordWrap(true);
ui->cfgs_grid->setItemDelegate(new GameGridDelegate(ui->cfgs_grid));
ui->cfgs_grid->setContextMenuPolicy(Qt::CustomContextMenu);
}
static void AddSaveDataDir(QStringList* dirs, QSet<QString>* seen, const QString& root,
const QString& title_id) {
const auto path = QDir(root).filePath(
@@ -252,7 +194,6 @@ ConfigurationListWidget::ConfigurationListWidget(QWidget* parent)
m_compatibility = new CompatibilityDatabase(
QCoreApplication::arguments().contains(QStringLiteral("--local")), this);
m_ui->setupUi(this);
ConfigureGameList(m_ui);
m_ui->refresh_action->setShortcuts(QKeySequence::Refresh);
m_ui->refresh_button->setDefaultAction(m_ui->refresh_action);
@@ -266,8 +207,9 @@ ConfigurationListWidget::ConfigurationListWidget(QWidget* parent)
m_ui->edit_button->setEnabled(false);
m_ui->cfgs_list->setContextMenuPolicy(Qt::CustomContextMenu);
m_ui->cfgs_list->setAllColumnsShowFocus(true);
m_ui->cfgs_list->setTextElideMode(Qt::ElideMiddle);
m_ui->cfgs_list->setIconSize(QSize(48, 48));
m_ui->cfgs_list->setRootIsDecorated(false);
m_ui->cfgs_list->setUniformRowHeights(true);
m_ui->cfgs_list->setSortingEnabled(true);
m_ui->cfgs_list->setColumnWidth(GAME_NAME_COLUMN, 320);
@@ -279,6 +221,8 @@ ConfigurationListWidget::ConfigurationListWidget(QWidget* parent)
m_ui->cfgs_list->setColumnWidth(GAME_STATUS_COLUMN, 150);
m_ui->cfgs_list->setColumnWidth(GAME_COMMENT_COLUMN, 240);
m_ui->cfgs_list->sortItems(GAME_NAME_COLUMN, Qt::AscendingOrder);
SetupGameGrid(m_ui->cfgs_list, m_ui->cfgs_grid);
m_ui->cfgs_grid->setContextMenuPolicy(Qt::CustomContextMenu);
connect(m_ui->refresh_action, &QAction::triggered, this,
&ConfigurationListWidget::ScanGameDirectory);
@@ -323,16 +267,12 @@ ConfigurationListWidget::ConfigurationListWidget(QWidget* parent)
connect(m_compatibility, &CompatibilityDatabase::Updated, this,
&ConfigurationListWidget::ApplyCompatibility);
m_ui->cfgs_list->setDragDropMode(QAbstractItemView::NoDragDrop);
ReadSettings();
m_ui->grid_size_slider->setVisible(m_ui->grid_view_button->isChecked());
const auto resize_grid = [this](int width) {
m_ui->cfgs_grid->setIconSize(QSize(width, width * 9 / 16));
};
resize_grid(m_ui->grid_size_slider->value());
connect(m_ui->grid_size_slider, &QSlider::valueChanged, this, [this, resize_grid](int width) {
resize_grid(width);
const int grid_size = m_ui->grid_size_slider->value();
m_ui->cfgs_grid->setIconSize(QSize(grid_size, grid_size));
connect(m_ui->grid_size_slider, &QSlider::valueChanged, this, [this](int width) {
m_ui->cfgs_grid->setIconSize(QSize(width, width));
if (!m_ui->grid_size_slider->isSliderDown()) {
WriteSettings();
}
@@ -1274,7 +1214,7 @@ void ConfigurationListWidget::SelectItem(QTreeWidgetItem* witem) {
auto* item = static_cast<ConfigurationItem*>(witem);
if (item == nullptr) {
m_selected_item = nullptr;
m_ui->cfgs_list->SetBackgroundImage({});
m_ui->game_views->SetBackgroundImage({});
m_ui->edit_button->setEnabled(false);
m_ui->delete_button->setEnabled(false);
emit Select();
@@ -1285,7 +1225,7 @@ void ConfigurationListWidget::SelectItem(QTreeWidgetItem* witem) {
m_ui->edit_button->setEnabled(!item->IsRunning());
m_selected_item = item;
m_ui->cfgs_list->SetBackgroundImage(item->GetInfo().basedir);
m_ui->game_views->SetBackgroundImage(item->GetInfo().basedir);
emit Select();
}
+1
View File
@@ -287,6 +287,7 @@ static QStringList CreateEmulatorArgs(const Configuration& info) {
AddFlag(args, info.skip_notice_screen, "--skip-notice-screen");
AddFlag(args, info.network_connected, "--network-connected");
AddFlag(args, info.tessellation_enabled, "--tessellation");
AddFlag(args, info.occlusion_queries_enabled, "--occlusion-queries");
args << "--vblank-frequency" << QString::number(info.vblank_frequency);
args << "--console-language" << QString::number(info.console_language);
AddFlag(args, info.vulkan_validation_enabled, "--vulkan-validation");
+6 -14
View File
@@ -1,32 +1,24 @@
/* The views get their transparent background from their own styleSheet in
configuration_list_widget.ui: Breeze paints black behind the scrollbars of views without one. */
QTreeWidget#cfgs_list,
QListView#cfgs_grid {
background: palette(base);
border: 1px solid palette(mid);
color: palette(text);
}
QListView#cfgs_grid::item {
padding: 8px;
}
GameListTreeWidget#cfgs_list {
background: transparent;
border: 1px solid palette(mid);
color: palette(text);
outline: 0;
}
GameListTreeWidget#cfgs_list::item {
QTreeWidget#cfgs_list::item {
background: transparent;
border-bottom: 1px solid palette(mid);
padding: 4px;
}
GameListTreeWidget#cfgs_list::item:selected {
QTreeWidget#cfgs_list::item:selected {
background: palette(highlight);
color: palette(highlighted-text);
}
GameListTreeWidget#cfgs_list QHeaderView::section {
QTreeWidget#cfgs_list QHeaderView::section {
background: palette(button);
color: palette(button-text);
border: 0;
+11 -26
View File
@@ -2027,7 +2027,7 @@ uint32_t* KYTY_SYSV_ABI AgcCbReleaseMem(CommandBuffer* buf, uint8_t action, uint
const uint32_t packet_action = static_cast<uint32_t>(action) & 0x3fu;
const uint32_t packet_event_index = (action >= 0x2fu ? 6u : 5u);
cmd[0] = KYTY_PM4(8, Pm4::IT_NOP, Pm4::R_RELEASE_MEM);
cmd[0] = KYTY_PM4(8, Pm4::IT_RELEASE_MEM, 0);
cmd[1] = packet_action | (packet_event_index << 8u) | ((packet_gcr_cntl & 0xfffu) << 12u) |
((static_cast<uint32_t>(cache_policy) & 0x3u) << 25u);
cmd[2] = ((static_cast<uint32_t>(dst) & 0x3u) << 16u) |
@@ -3510,13 +3510,7 @@ int KYTY_SYSV_ABI AgcSetPacketPredication(uint32_t* packet, uint32_t predication
int KYTY_SYSV_ABI AgcSetNop(uint32_t* packet) {
PRINT_NAME();
const auto size = AgcGetPacketSize(packet);
packet[0] = (packet[0] & 0xffff0003u) | (Pm4::IT_NOP << 8u);
// Prevent the disabled payload from being treated as an executable marker.
if (size > 1 && (packet[1] & 0xffff0000u) == 0x68750000u) {
packet[1] &= 0xffffu;
}
return OK;
}
@@ -3778,27 +3772,21 @@ int KYTY_SYSV_ABI AgcQueueEndOfPipeActionPatchAddress(uint32_t* cmd,
const volatile Label* address) {
PRINT_NAME();
// Not sure
AgcTrace("\t cmd = 0x%016" PRIx64 "\n"
"\t address = 0x%016" PRIx64 "\n",
reinterpret_cast<uint64_t>(cmd), reinterpret_cast<uint64_t>(address));
EXIT_NOT_IMPLEMENTED(cmd == nullptr);
auto vaddr = reinterpret_cast<uint64_t>(address);
auto op = (cmd[0] >> 8u) & 0xffu;
if (((cmd[0] >> 8u) & 0xffu) != Pm4::IT_RELEASE_MEM) {
return GRAPHICS5_ERROR_INVALID_PACKET;
}
if ((op == Pm4::IT_NOP && KYTY_PM4_R(cmd[0]) == Pm4::R_RELEASE_MEM) ||
op == Pm4::IT_RELEASE_MEM) {
cmd[3] = static_cast<uint32_t>(vaddr & 0xffffffffu);
cmd[4] = static_cast<uint32_t>((vaddr >> 32u) & 0xffffffffu);
} else if (op == Pm4::IT_EVENT_WRITE_EOP) {
cmd[2] = static_cast<uint32_t>(vaddr & 0xffffffffu);
cmd[3] = (cmd[3] & 0xffff0000u) | static_cast<uint32_t>((vaddr >> 32u) & 0xffffu);
} else {
EXIT("unsupported queueEndOfPipeAction packet for address patch: 0x%08" PRIx32 "\n",
cmd[0]);
// A kSendIntCtx release has no address.
if (((cmd[2] >> 24u) & 0x7u) != 4u) {
const auto vaddr = reinterpret_cast<uint64_t>(address);
cmd[3] = (cmd[3] & 0x3u) | (static_cast<uint32_t>(vaddr) & ~0x3u);
cmd[4] = static_cast<uint32_t>(vaddr >> 32u);
}
return OK;
@@ -3813,10 +3801,7 @@ int KYTY_SYSV_ABI AgcQueueEndOfPipeActionPatchData(uint32_t* cmd, uint64_t data)
EXIT_NOT_IMPLEMENTED(cmd == nullptr);
const auto op = (cmd[0] >> 8u) & 0xffu;
const bool is_release_mem = op == Pm4::IT_RELEASE_MEM ||
(op == Pm4::IT_NOP && KYTY_PM4_R(cmd[0]) == Pm4::R_RELEASE_MEM);
if (!is_release_mem) {
if (((cmd[0] >> 8u) & 0xffu) != Pm4::IT_RELEASE_MEM) {
return GRAPHICS5_ERROR_INVALID_PACKET;
}
@@ -4115,7 +4100,7 @@ static TessellationDriverState g_tessellation_driver_state {};
static void submit_dcb(uint32_t* dcb, uint32_t size_in_dwords) {
GraphicsDbgDumpDcb("d", size_in_dwords, dcb);
EXIT_IF(g_renderer == nullptr);
g_renderer->GetGpu().Submit(std::span {dcb, size_in_dwords}, {});
g_renderer->GetGpu().Submit(std::span {dcb, size_in_dwords});
}
int KYTY_SYSV_ABI AgcDriverSubmitDcb(const Packet* packet) {
+9
View File
@@ -16,6 +16,14 @@ static int KYTY_SYSV_ABI DeviceServiceInitialize() {
return OK;
}
static int KYTY_SYSV_ABI DeviceServiceGetEventState(int32_t device_class) {
PRINT_NAME();
LOGF("\t device_class = %" PRId32 "\n", device_class);
return 0;
}
static int KYTY_SYSV_ABI DeviceServiceQueryDeviceInfo(uint32_t device_handle, void* out_info) {
PRINT_NAME();
@@ -30,6 +38,7 @@ static int KYTY_SYSV_ABI DeviceServiceQueryDeviceInfo(uint32_t device_handle, vo
LIB_DEFINE(InitDeviceService_1) {
LIB_FUNC("84fDxStrG44", DeviceService::DeviceServiceInitialize);
LIB_FUNC("9ddRUOV8Q5A", DeviceService::DeviceServiceGetEventState);
LIB_FUNC("UNMEa+5lrUA", DeviceService::DeviceServiceQueryDeviceInfo);
}
+402
View File
@@ -0,0 +1,402 @@
#include "common/abi.h"
#include "common/logging/log.h"
#include "libs/libs.h"
#include "loader/symbolDatabase.h"
#include <algorithm>
#include <cstddef>
#include <cstdlib>
// The emulator's stb_image implementation. Guest JPEG data reaches it here, and zeroed allocations
// keep a scan that uses a Huffman table no DHT defined (an MJPEG frame) from reading uninitialised
// tables.
#define STB_IMAGE_IMPLEMENTATION
#define STBI_NO_SIMD
#define STBI_MALLOC(size) calloc(1, size)
#define STBI_REALLOC(pointer, size) realloc(pointer, size)
#define STBI_FREE(pointer) free(pointer)
#include "stb_image.h"
namespace Libs {
LIB_VERSION("JpegDec", 1, "JpegDec", 1, 1);
namespace JpegDec {
constexpr int32_t JPEG_DEC_ERROR_INVALID_ADDR = static_cast<int32_t>(0x80650001u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_SIZE = static_cast<int32_t>(0x80650002u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_PARAM = static_cast<int32_t>(0x80650003u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_HANDLE = static_cast<int32_t>(0x80650004u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_COEF_MEMORY = static_cast<int32_t>(0x80650005u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_WORK_MEMORY = static_cast<int32_t>(0x80650006u);
constexpr int32_t JPEG_DEC_ERROR_INVALID_DATA = static_cast<int32_t>(0x80650010u);
constexpr int32_t JPEG_DEC_ERROR_UNSUPPORT_DATA = static_cast<int32_t>(0x80650011u);
constexpr int32_t JPEG_DEC_ERROR_DECODE_ERROR = static_cast<int32_t>(0x80650012u);
constexpr uint32_t JPEG_DEC_ATTRIBUTE_CSC_STD_SAMPLING = 1;
constexpr uint32_t JPEG_DEC_ATTRIBUTE_CSC_FULL_SAMPLING = 2;
constexpr uint16_t JPEG_DEC_DECODE_MODE_MJPEG = 1;
constexpr uint16_t JPEG_DEC_PIXEL_FORMAT_R8G8B8A8 = 0;
constexpr uint16_t JPEG_DEC_PIXEL_FORMAT_B8G8R8A8 = 1;
constexpr uint16_t JPEG_DEC_PIXEL_FORMAT_Y8U8Y8V8 = 10;
constexpr uint16_t JPEG_DEC_PIXEL_FORMAT_Y8 = 11;
constexpr uint16_t JPEG_DEC_COLOR_SPACE_UNKNOWN = 0;
constexpr uint16_t JPEG_DEC_COLOR_SPACE_YCC = 1;
constexpr uint16_t JPEG_DEC_COLOR_SPACE_GRAYSCALE = 2;
struct JpegDecCreateParam {
uint32_t this_size;
uint32_t attribute;
uint32_t max_image_width;
};
struct JpegDecDecodeParam {
const void* jpeg_mem_addr;
void* image_mem_addr;
void* coefficient_mem_addr;
uint32_t jpeg_mem_size;
uint32_t image_mem_size;
uint32_t coefficient_mem_size;
uint16_t decode_mode;
uint16_t down_scale;
uint16_t pixel_format;
uint16_t alpha_value;
uint32_t image_pitch;
};
struct JpegDecParseParam {
const void* jpeg_mem_addr;
uint32_t jpeg_mem_size;
uint16_t decode_mode;
uint16_t down_scale;
};
struct JpegDecImageInfo {
uint32_t image_width;
uint32_t image_height;
uint16_t color_space;
uint16_t component_number;
uint8_t component_id[4];
uint8_t sampling_factor[4];
uint32_t coefficient_mem_size;
uint32_t suitable_csc_attribute;
uint32_t output_image_width;
uint32_t output_image_height;
};
static_assert(sizeof(JpegDecCreateParam) == 12 && sizeof(JpegDecDecodeParam) == 48);
static_assert(sizeof(JpegDecParseParam) == 16 && sizeof(JpegDecImageInfo) == 36);
struct JpegDecContext {
uint64_t magic;
uint32_t attribute;
uint32_t max_image_width;
};
constexpr uint64_t JPEG_DEC_CONTEXT_MAGIC = 0x4b5954594a504744ull; // KYTYJPGD
static int32_t CheckCreateParam(const JpegDecCreateParam* param) {
if (param == nullptr) {
return JPEG_DEC_ERROR_INVALID_ADDR;
}
if (param->this_size != sizeof(JpegDecCreateParam)) {
return JPEG_DEC_ERROR_INVALID_SIZE;
}
if ((param->attribute != JPEG_DEC_ATTRIBUTE_CSC_STD_SAMPLING &&
param->attribute != JPEG_DEC_ATTRIBUTE_CSC_FULL_SAMPLING) ||
param->max_image_width == 0 || param->max_image_width > 0xffffu) {
return JPEG_DEC_ERROR_INVALID_PARAM;
}
return 0;
}
static bool CheckScale(uint16_t decode_mode, uint16_t down_scale) {
return decode_mode <= JPEG_DEC_DECODE_MODE_MJPEG &&
(down_scale == 1 || down_scale == 2 || down_scale == 4 || down_scale == 8);
}
// Fills the image information from the frame header (T.81 B.2.2) that precedes the first scan.
static int32_t ParseHeader(const void* jpeg, uint32_t size, uint16_t down_scale,
JpegDecImageInfo* info) {
const auto* data = static_cast<const uint8_t*>(jpeg);
if (size < 4 || data[0] != 0xff || data[1] != 0xd8) {
return JPEG_DEC_ERROR_INVALID_DATA;
}
for (uint32_t offset = 2; offset + 4 <= size;) {
const uint8_t marker = data[offset + 1];
if (data[offset] != 0xff || marker == 0xda || marker == 0xd9) {
return JPEG_DEC_ERROR_INVALID_DATA;
}
if (marker == 0xff) {
offset++;
continue;
}
const uint32_t length = (data[offset + 2] << 8u) | data[offset + 3];
const auto* segment = data + offset + 4;
if (length < 2 || offset + 2 + length > size) {
return JPEG_DEC_ERROR_INVALID_DATA;
}
offset += 2 + length;
// SOF0-SOF15 except DHT, JPG and DAC.
if (marker < 0xc0 || marker > 0xcf || marker == 0xc4 || marker == 0xc8 || marker == 0xcc) {
continue;
}
const uint32_t components = length < 8 ? 0 : segment[5];
if (components == 0 || length < 8 + components * 3 || (segment[3] | segment[4]) == 0) {
return JPEG_DEC_ERROR_INVALID_DATA;
}
*info = {};
info->image_height = (segment[1] << 8u) | segment[2];
info->image_width = (segment[3] << 8u) | segment[4];
// Baseline, extended and progressive Huffman 8-bit frames with a known height.
if (marker > 0xc2 || segment[0] != 8 || components > 4 || info->image_height == 0) {
return JPEG_DEC_ERROR_UNSUPPORT_DATA;
}
uint32_t h_max = 1;
uint32_t v_max = 1;
for (uint32_t i = 0; i < components; i++) {
const uint8_t factor = segment[7 + i * 3];
if ((factor >> 4u) - 1u > 3u || (factor & 0xfu) - 1u > 3u) {
return JPEG_DEC_ERROR_INVALID_DATA;
}
info->component_id[i] = segment[6 + i * 3];
info->sampling_factor[i] = factor;
h_max = std::max<uint32_t>(h_max, factor >> 4u);
v_max = std::max<uint32_t>(v_max, factor & 0xfu);
}
info->component_number = components;
if (marker == 0xc2) {
// Progressive scans keep every 16-bit DCT coefficient of the MCU-padded image.
const uint32_t mcu_x = (info->image_width + h_max * 8 - 1) / (h_max * 8);
const uint32_t mcu_y = (info->image_height + v_max * 8 - 1) / (v_max * 8);
for (uint32_t i = 0; i < components; i++) {
info->coefficient_mem_size += mcu_x * (info->sampling_factor[i] >> 4u) * mcu_y *
(info->sampling_factor[i] & 0xfu) * 128u;
}
}
switch (components) {
case 1: info->color_space = JPEG_DEC_COLOR_SPACE_GRAYSCALE; break;
case 3: info->color_space = JPEG_DEC_COLOR_SPACE_YCC; break;
default: info->color_space = JPEG_DEC_COLOR_SPACE_UNKNOWN; break;
}
// Standard sampling: grayscale, or 1x1 chroma with luma at most 2x2 (4:4:4/4:2:2/4:2:0).
const bool standard =
components == 1 || (info->sampling_factor[1] == 0x11 &&
info->sampling_factor[2] == 0x11 && h_max <= 2 && v_max <= 2);
info->suitable_csc_attribute =
standard ? JPEG_DEC_ATTRIBUTE_CSC_STD_SAMPLING : JPEG_DEC_ATTRIBUTE_CSC_FULL_SAMPLING;
info->output_image_width = (info->image_width + down_scale - 1) / down_scale;
info->output_image_height = (info->image_height + down_scale - 1) / down_scale;
return 0;
}
return JPEG_DEC_ERROR_INVALID_DATA;
}
static uint8_t Luma(const uint8_t* p) {
return static_cast<uint8_t>((77 * p[0] + 150 * p[1] + 29 * p[2] + 128) >> 8);
}
// JFIF chroma of the average of two pixels, rounding half down so a saturated pair stays 255.
static uint8_t Chroma(const uint8_t* p, const uint8_t* q, int r, int g, int b) {
return static_cast<uint8_t>(
(r * (p[0] + q[0]) + g * (p[1] + q[1]) + b * (p[2] + q[2]) + 65791) >> 9);
}
static int32_t KYTY_SYSV_ABI JpegDecQueryMemorySize(const JpegDecCreateParam* param) {
PRINT_NAME();
const auto result = CheckCreateParam(param);
return result != 0 ? result : static_cast<int32_t>(sizeof(JpegDecContext));
}
static int32_t KYTY_SYSV_ABI JpegDecCreate(const JpegDecCreateParam* param, void* memory_address,
uint32_t memory_size, void** handle) {
PRINT_NAME();
if (const auto result = CheckCreateParam(param); result != 0) {
return result;
}
if (memory_address == nullptr || handle == nullptr) {
return JPEG_DEC_ERROR_INVALID_ADDR;
}
if (memory_size < sizeof(JpegDecContext)) {
return JPEG_DEC_ERROR_INVALID_SIZE;
}
*static_cast<JpegDecContext*>(memory_address) = {JPEG_DEC_CONTEXT_MAGIC, param->attribute,
param->max_image_width};
*handle = memory_address;
return 0;
}
static int32_t KYTY_SYSV_ABI JpegDecDelete(void* handle) {
PRINT_NAME();
auto* ctx = static_cast<JpegDecContext*>(handle);
if (ctx == nullptr || ctx->magic != JPEG_DEC_CONTEXT_MAGIC) {
return JPEG_DEC_ERROR_INVALID_HANDLE;
}
ctx->magic = 0;
return 0;
}
static int32_t KYTY_SYSV_ABI JpegDecParseHeader(const JpegDecParseParam* param,
JpegDecImageInfo* image_info) {
PRINT_NAME();
if (param == nullptr || image_info == nullptr || param->jpeg_mem_addr == nullptr) {
return JPEG_DEC_ERROR_INVALID_ADDR;
}
if (param->jpeg_mem_size == 0) {
return JPEG_DEC_ERROR_INVALID_SIZE;
}
if (!CheckScale(param->decode_mode, param->down_scale)) {
return JPEG_DEC_ERROR_INVALID_PARAM;
}
return ParseHeader(param->jpeg_mem_addr, param->jpeg_mem_size, param->down_scale, image_info);
}
static int32_t KYTY_SYSV_ABI JpegDecDecode(void* handle, const JpegDecDecodeParam* param,
JpegDecImageInfo* image_info) {
PRINT_NAME();
const auto* ctx = static_cast<const JpegDecContext*>(handle);
if (ctx == nullptr || ctx->magic != JPEG_DEC_CONTEXT_MAGIC) {
return JPEG_DEC_ERROR_INVALID_HANDLE;
}
if (param == nullptr || param->jpeg_mem_addr == nullptr || param->image_mem_addr == nullptr ||
(param->pixel_format != JPEG_DEC_PIXEL_FORMAT_Y8 &&
(reinterpret_cast<uintptr_t>(param->image_mem_addr) & 3u) != 0)) {
return JPEG_DEC_ERROR_INVALID_ADDR;
}
const auto format = param->pixel_format;
LOGF("\t jpeg_mem_size = %" PRIu32 ", image_mem_size = %" PRIu32 ", mode = %" PRIu16
", down_scale = %" PRIu16 ", pixel_format = %" PRIu16 ", image_pitch = %" PRIu32 "\n",
param->jpeg_mem_size, param->image_mem_size, param->decode_mode, param->down_scale, format,
param->image_pitch);
if (param->jpeg_mem_size == 0 || param->image_mem_size == 0) {
return JPEG_DEC_ERROR_INVALID_SIZE;
}
if (!CheckScale(param->decode_mode, param->down_scale) || param->alpha_value > 255 ||
(format != JPEG_DEC_PIXEL_FORMAT_R8G8B8A8 && format != JPEG_DEC_PIXEL_FORMAT_B8G8R8A8 &&
format != JPEG_DEC_PIXEL_FORMAT_Y8U8Y8V8 && format != JPEG_DEC_PIXEL_FORMAT_Y8) ||
(format != JPEG_DEC_PIXEL_FORMAT_Y8 && (param->image_pitch & 3u) != 0)) {
return JPEG_DEC_ERROR_INVALID_PARAM;
}
JpegDecImageInfo info {};
if (const auto result =
ParseHeader(param->jpeg_mem_addr, param->jpeg_mem_size, param->down_scale, &info);
result != 0) {
return result;
}
if (info.color_space == JPEG_DEC_COLOR_SPACE_UNKNOWN) {
return JPEG_DEC_ERROR_UNSUPPORT_DATA;
}
if (info.output_image_width > ctx->max_image_width ||
info.suitable_csc_attribute > ctx->attribute) {
return JPEG_DEC_ERROR_INVALID_WORK_MEMORY;
}
if (info.coefficient_mem_size > param->coefficient_mem_size ||
(info.coefficient_mem_size != 0 && param->coefficient_mem_addr == nullptr)) {
return JPEG_DEC_ERROR_INVALID_COEF_MEMORY;
}
const uint32_t width = info.output_image_width;
const uint32_t height = info.output_image_height;
const bool large = width > 32767 || height > 32767;
if (large && image_info == nullptr) {
return JPEG_DEC_ERROR_INVALID_ADDR;
}
uint32_t row = width * 4;
if (format == JPEG_DEC_PIXEL_FORMAT_Y8U8Y8V8) {
row = (width + 1) / 2 * 4;
} else if (format == JPEG_DEC_PIXEL_FORMAT_Y8) {
row = width;
}
const uint32_t pitch = param->image_pitch == 0 ? row : param->image_pitch;
if (pitch < row) {
return JPEG_DEC_ERROR_INVALID_PARAM;
}
if (static_cast<uint64_t>(pitch) * height > param->image_mem_size) {
return JPEG_DEC_ERROR_INVALID_SIZE;
}
int src_width = 0;
int src_height = 0;
auto* rgb = stbi_load_from_memory(static_cast<const stbi_uc*>(param->jpeg_mem_addr),
static_cast<int>(param->jpeg_mem_size), &src_width,
&src_height, nullptr, 3);
if (rgb == nullptr || static_cast<uint32_t>(src_width) != info.image_width ||
static_cast<uint32_t>(src_height) != info.image_height) {
stbi_image_free(rgb);
return JPEG_DEC_ERROR_DECODE_ERROR;
}
// Each output pixel averages its down_scale x down_scale source block, packed in place: an
// output pixel never lands past a source pixel still to be read.
const uint32_t scale = param->down_scale;
if (scale > 1) {
auto* out = rgb;
for (uint32_t y = 0; y < height; y++) {
for (uint32_t x = 0; x < width; x++) {
const uint32_t x_end = std::min((x + 1) * scale, info.image_width);
const uint32_t y_end = std::min((y + 1) * scale, info.image_height);
const uint32_t count = (x_end - x * scale) * (y_end - y * scale);
uint32_t sum[3] = {};
for (uint32_t sy = y * scale; sy < y_end; sy++) {
for (uint32_t sx = x * scale; sx < x_end; sx++) {
const auto* p = rgb + (static_cast<size_t>(sy) * info.image_width + sx) * 3;
sum[0] += p[0];
sum[1] += p[1];
sum[2] += p[2];
}
}
for (const uint32_t c: sum) {
*out++ = static_cast<uint8_t>((c + count / 2) / count);
}
}
}
}
const auto alpha = static_cast<uint8_t>(param->alpha_value);
const auto* p = rgb;
for (uint32_t y = 0; y < height; y++) {
auto* dst = static_cast<uint8_t*>(param->image_mem_addr) + static_cast<size_t>(y) * pitch;
for (uint32_t x = 0; x < width; x++, p += 3) {
switch (format) {
case JPEG_DEC_PIXEL_FORMAT_R8G8B8A8:
dst = std::copy_n(p, 3, dst);
*dst++ = alpha;
break;
case JPEG_DEC_PIXEL_FORMAT_B8G8R8A8:
*dst++ = p[2];
*dst++ = p[1];
*dst++ = p[0];
*dst++ = alpha;
break;
case JPEG_DEC_PIXEL_FORMAT_Y8: *dst++ = Luma(p); break;
default:
if (x % 2 == 0) {
// An odd width repeats the last pixel.
const auto* q = x + 1 < width ? p + 3 : p;
*dst++ = Luma(p);
*dst++ = Chroma(p, q, -43, -85, 128);
*dst++ = Luma(q);
*dst++ = Chroma(p, q, 128, -107, -21);
}
break;
}
}
}
stbi_image_free(rgb);
if (image_info != nullptr) {
*image_info = info;
}
return large ? 0 : static_cast<int32_t>((width << 16u) | height);
}
} // namespace JpegDec
LIB_DEFINE(InitJpegDec_1) {
PRINT_NAME_ENABLE(true);
LIB_FUNC("uNAUmANZMEw", JpegDec::JpegDecQueryMemorySize);
LIB_FUNC("JPh3Zgg0Zwc", JpegDec::JpegDecCreate);
LIB_FUNC("1kzQRoWEgSA", JpegDec::JpegDecDecode);
LIB_FUNC("Hwh11+m5KoI", JpegDec::JpegDecDelete);
LIB_FUNC("LSinoSQH790", JpegDec::JpegDecParseHeader);
}
} // namespace Libs
+240
View File
@@ -0,0 +1,240 @@
#include "common/abi.h"
#include "common/logging/log.h"
#include "libs/libs.h"
#include "loader/symbolDatabase.h"
#include <algorithm>
#include <cstddef>
#include <cstring>
#include <vector>
#define STB_IMAGE_WRITE_IMPLEMENTATION
#define STB_IMAGE_WRITE_STATIC
#define STBIWDEF static inline
#define STBI_WRITE_NO_STDIO
#include "stb_image_write.h"
namespace Libs {
LIB_VERSION("JpegEnc", 1, "JpegEnc", 1, 1);
namespace JpegEnc {
constexpr int32_t JPEG_ENC_ERROR_INVALID_ADDR = static_cast<int32_t>(0x80650101u);
constexpr int32_t JPEG_ENC_ERROR_INVALID_SIZE = static_cast<int32_t>(0x80650102u);
constexpr int32_t JPEG_ENC_ERROR_INVALID_PARAM = static_cast<int32_t>(0x80650103u);
constexpr int32_t JPEG_ENC_ERROR_INVALID_HANDLE = static_cast<int32_t>(0x80650104u);
constexpr int32_t JPEG_ENC_ERROR_DATA_OVERFLOW = static_cast<int32_t>(0x80650110u);
constexpr uint16_t JPEG_ENC_PIXEL_FORMAT_R8G8B8A8 = 0;
constexpr uint16_t JPEG_ENC_PIXEL_FORMAT_B8G8R8A8 = 1;
constexpr uint16_t JPEG_ENC_PIXEL_FORMAT_Y8U8Y8V8 = 10;
constexpr uint16_t JPEG_ENC_PIXEL_FORMAT_Y8 = 11;
constexpr uint16_t JPEG_ENC_ENCODE_MODE_MJPEG = 1;
constexpr uint16_t JPEG_ENC_COLOR_SPACE_YCC = 1;
constexpr uint16_t JPEG_ENC_COLOR_SPACE_GRAYSCALE = 2;
constexpr uint8_t JPEG_ENC_SAMPLING_TYPE_FULL = 0;
constexpr uint8_t JPEG_ENC_SAMPLING_TYPE_422 = 1;
constexpr uint8_t JPEG_ENC_SAMPLING_TYPE_420 = 2;
struct JpegEncCreateParam {
uint32_t this_size;
uint32_t attribute;
};
struct JpegEncEncodeParam {
const uint8_t* image_mem_addr;
uint8_t* jpeg_mem_addr;
uint32_t image_mem_size;
uint32_t jpeg_mem_size;
uint32_t image_width;
uint32_t image_height;
uint32_t image_pitch;
uint16_t pixel_format;
uint16_t encode_mode;
uint16_t color_space;
uint8_t sampling_type;
uint8_t compression_ratio;
int32_t restart_interval;
};
struct JpegEncOutputInfo {
uint32_t data_size;
uint32_t processed_height;
};
static_assert(sizeof(JpegEncCreateParam) == 8 && sizeof(JpegEncEncodeParam) == 48);
struct JpegEncContext {
uint64_t magic;
};
constexpr uint64_t JPEG_ENC_CONTEXT_MAGIC = 0x4b5954594a504745ull; // KYTYJPGE
struct JpegEncWriter {
uint8_t* dst;
uint32_t capacity;
uint32_t size;
bool overflow;
};
static int32_t CheckCreateParam(const JpegEncCreateParam* param) {
if (param == nullptr) {
return JPEG_ENC_ERROR_INVALID_ADDR;
}
if (param->this_size != sizeof(JpegEncCreateParam)) {
return JPEG_ENC_ERROR_INVALID_SIZE;
}
return param->attribute != 0 ? JPEG_ENC_ERROR_INVALID_PARAM : 0;
}
static void JpegEncWrite(void* context, void* data, int size) {
auto* writer = static_cast<JpegEncWriter*>(context);
if (writer->overflow || static_cast<uint32_t>(size) > writer->capacity - writer->size) {
writer->overflow = true;
return;
}
std::memcpy(writer->dst + writer->size, data, size);
writer->size += size;
}
static int32_t KYTY_SYSV_ABI JpegEncQueryMemorySize(const JpegEncCreateParam* param) {
PRINT_NAME();
const auto result = CheckCreateParam(param);
return result != 0 ? result : static_cast<int32_t>(sizeof(JpegEncContext));
}
static int32_t KYTY_SYSV_ABI JpegEncCreate(const JpegEncCreateParam* param, void* memory_address,
uint32_t memory_size, void** handle) {
PRINT_NAME();
if (const auto result = CheckCreateParam(param); result != 0) {
return result;
}
if (memory_address == nullptr || handle == nullptr) {
return JPEG_ENC_ERROR_INVALID_ADDR;
}
if (memory_size < sizeof(JpegEncContext)) {
return JPEG_ENC_ERROR_INVALID_SIZE;
}
static_cast<JpegEncContext*>(memory_address)->magic = JPEG_ENC_CONTEXT_MAGIC;
*handle = memory_address;
return 0;
}
static int32_t KYTY_SYSV_ABI JpegEncDelete(void* handle) {
PRINT_NAME();
auto* ctx = static_cast<JpegEncContext*>(handle);
if (ctx == nullptr || ctx->magic != JPEG_ENC_CONTEXT_MAGIC) {
return JPEG_ENC_ERROR_INVALID_HANDLE;
}
ctx->magic = 0;
return 0;
}
// stb_image_write picks the chroma sampling from the quality, always writes three components and
// full headers, and writes no restart markers; these only change the bytes of a valid JPEG.
static int32_t KYTY_SYSV_ABI JpegEncEncode(void* handle, const JpegEncEncodeParam* param,
JpegEncOutputInfo* output_info) {
PRINT_NAME();
const auto* ctx = static_cast<const JpegEncContext*>(handle);
if (ctx == nullptr || ctx->magic != JPEG_ENC_CONTEXT_MAGIC) {
return JPEG_ENC_ERROR_INVALID_HANDLE;
}
if (param == nullptr || param->image_mem_addr == nullptr || param->jpeg_mem_addr == nullptr ||
(param->pixel_format != JPEG_ENC_PIXEL_FORMAT_Y8 &&
(reinterpret_cast<uintptr_t>(param->image_mem_addr) & 3u) != 0)) {
return JPEG_ENC_ERROR_INVALID_ADDR;
}
const auto format = param->pixel_format;
const uint32_t width = param->image_width;
const uint32_t height = param->image_height;
LOGF("\t %" PRIu32 "x%" PRIu32 ", pitch = %" PRIu32 ", pixel_format = %" PRIu16
", mode = %" PRIu16 ", color_space = %" PRIu16 ", sampling = %" PRIu8 ", ratio = %" PRIu8
", restart = %" PRId32 "\n",
width, height, param->image_pitch, format, param->encode_mode, param->color_space,
param->sampling_type, param->compression_ratio, param->restart_interval);
const bool gray = param->color_space == JPEG_ENC_COLOR_SPACE_GRAYSCALE;
uint32_t row = width * 4;
if (format == JPEG_ENC_PIXEL_FORMAT_Y8U8Y8V8) {
row = (width + 1) / 2 * 4;
} else if (format == JPEG_ENC_PIXEL_FORMAT_Y8) {
row = width;
}
// Grayscale takes full sampling; YCbCr takes 4:2:2 or 4:2:0 from a color format.
const auto sampling = param->sampling_type;
const bool ycc =
param->color_space == JPEG_ENC_COLOR_SPACE_YCC && format != JPEG_ENC_PIXEL_FORMAT_Y8 &&
(sampling == JPEG_ENC_SAMPLING_TYPE_422 || sampling == JPEG_ENC_SAMPLING_TYPE_420);
if (width == 0 || width > 0xffff || height == 0 || height > 0xffff ||
param->image_pitch < row ||
(format != JPEG_ENC_PIXEL_FORMAT_Y8 && (param->image_pitch & 3u) != 0) ||
(format != JPEG_ENC_PIXEL_FORMAT_R8G8B8A8 && format != JPEG_ENC_PIXEL_FORMAT_B8G8R8A8 &&
format != JPEG_ENC_PIXEL_FORMAT_Y8U8Y8V8 && format != JPEG_ENC_PIXEL_FORMAT_Y8) ||
param->encode_mode > JPEG_ENC_ENCODE_MODE_MJPEG ||
!(ycc || (gray && sampling == JPEG_ENC_SAMPLING_TYPE_FULL)) ||
param->restart_interval < -1 || param->restart_interval > 0xffff) {
return JPEG_ENC_ERROR_INVALID_PARAM;
}
if (param->jpeg_mem_size == 0 ||
static_cast<uint64_t>(param->image_pitch) * height > param->image_mem_size) {
return JPEG_ENC_ERROR_INVALID_SIZE;
}
const uint32_t components = gray ? 1 : 3;
std::vector<uint8_t> pixels(static_cast<size_t>(width) * height * components);
auto* dst = pixels.data();
for (uint32_t y = 0; y < height; y++) {
const auto* src = param->image_mem_addr + static_cast<size_t>(y) * param->image_pitch;
for (uint32_t x = 0; x < width; x++) {
int rgb[3] = {src[x], src[x], src[x]};
if (format == JPEG_ENC_PIXEL_FORMAT_Y8U8Y8V8) {
// JFIF YCbCr to RGB from the YUYV pair that holds this pixel; gray keeps the luma.
const auto* pair = src + x / 2 * 4;
const int luma = pair[x % 2 * 2] * 256 + 128;
const int cb = gray ? 0 : pair[1] - 128;
const int cr = gray ? 0 : pair[3] - 128;
rgb[0] = std::clamp((luma + 359 * cr) >> 8, 0, 255);
rgb[1] = std::clamp((luma - 88 * cb - 183 * cr) >> 8, 0, 255);
rgb[2] = std::clamp((luma + 454 * cb) >> 8, 0, 255);
} else if (format != JPEG_ENC_PIXEL_FORMAT_Y8) {
const auto* p = src + x * 4;
const bool bgr = format == JPEG_ENC_PIXEL_FORMAT_B8G8R8A8;
rgb[0] = p[bgr ? 2 : 0];
rgb[1] = p[1];
rgb[2] = p[bgr ? 0 : 2];
}
if (gray) {
*dst++ =
static_cast<uint8_t>((77 * rgb[0] + 150 * rgb[1] + 29 * rgb[2] + 128) >> 8);
continue;
}
for (const int c: rgb) {
*dst++ = static_cast<uint8_t>(c);
}
}
}
// A lower compression ratio means higher quality: 0 maps to stb quality 100 and 255 to 1.
JpegEncWriter writer {param->jpeg_mem_addr, param->jpeg_mem_size, 0, false};
stbi_write_jpg_to_func(JpegEncWrite, &writer, static_cast<int>(width), static_cast<int>(height),
static_cast<int>(components), pixels.data(),
100 - param->compression_ratio * 99 / 255);
if (output_info != nullptr) {
output_info->data_size = writer.size;
output_info->processed_height = writer.overflow ? 0 : height;
}
return writer.overflow ? JPEG_ENC_ERROR_DATA_OVERFLOW : static_cast<int32_t>(writer.size);
}
} // namespace JpegEnc
LIB_DEFINE(InitJpegEnc_1) {
PRINT_NAME_ENABLE(true);
LIB_FUNC("o6ZgXfFdWXQ", JpegEnc::JpegEncQueryMemorySize);
LIB_FUNC("K+rocojkr-I", JpegEnc::JpegEncCreate);
LIB_FUNC("QbrU0cUghEM", JpegEnc::JpegEncEncode);
LIB_FUNC("j1LyMdaM+C0", JpegEnc::JpegEncDelete);
}
} // namespace Libs
-3
View File
@@ -36,9 +36,6 @@ static int KYTY_SYSV_ABI UsbdHandleEventsTimeout() {
static int64_t KYTY_SYSV_ABI UsbdGetDeviceList(void*** list) {
PRINT_NAME();
static void* empty_list[] = {nullptr};
*list = empty_list;
return 0;
}
+4
View File
@@ -92,6 +92,8 @@ LIB_DEFINE(InitFont_1);
LIB_DEFINE(InitFontFt_1);
LIB_DEFINE(InitAgcDriver_1);
LIB_DEFINE(InitHmd2_1);
LIB_DEFINE(InitJpegDec_1);
LIB_DEFINE(InitJpegEnc_1);
LIB_DEFINE(InitVrTracker2_1);
LIB_DEFINE(InitLibKernel_1);
LIB_DEFINE(InitNet_1);
@@ -132,6 +134,8 @@ void InitAll(Loader::SymbolDatabase* s) {
InitDeviceService_1(s);
InitUsbd_1(s);
InitHmd2_1(s);
InitJpegDec_1(s);
InitJpegEnc_1(s);
InitVrTracker2_1(s);
InitLibKernel_1(s);
LibMouse::InitMouse_1(s);
+1 -1
View File
@@ -466,7 +466,7 @@ Network::Id Network::HttpCreateTemplate(Id http_ctx_id, const char* user_agent,
HttpTemplate tn {};
tn.used = true;
tn.http_ver = http_ver;
tn.user_agent = std::string(user_agent);
tn.user_agent = user_agent != nullptr ? user_agent : "";
tn.is_auto_proxy_conf = is_auto_proxy_conf;
tn.http_ctx_id = http_ctx_id.GetId();
tn.nonblock = false;
+2
View File
@@ -78,6 +78,7 @@ static void PrintUsage() {
" Implies --vulkan-validation; very slow.\n");
::printf(" --shader-validation Enable shader validation.\n");
::printf(" --tessellation Draw tessellation patches; skipped by default.\n");
::printf(" --occlusion-queries Count occlusion query samples on the GPU.\n");
::printf(" --shader-optimization-type <value> None, Size, or Performance.\n");
::printf(" --shader-log-direction <value> Silent, Console, or File.\n");
::printf(" --shader-log-folder <path> Shader log output folder.\n");
@@ -202,6 +203,7 @@ static constexpr Flag FLAGS[] = {
{"--gpu-assisted-validation", &ConfigOptions::gpu_assisted_validation_enabled},
{"--shader-validation", &ConfigOptions::shader_validation_enabled},
{"--tessellation", &ConfigOptions::tessellation_enabled},
{"--occlusion-queries", &ConfigOptions::occlusion_queries_enabled},
{"--command-buffer-dump", &ConfigOptions::command_buffer_dump_enabled},
{"--graphics-debug-dump", &ConfigOptions::graphics_debug_dump_enabled},
{"--profile", &ConfigOptions::profiler_enabled},
+13
View File
@@ -1,4 +1,5 @@
#include "common/abi.h"
#include "libs/network.h"
#include "loader/symbolDatabase.h"
#include <array>
@@ -162,6 +163,17 @@ void TestBuildHonoursOption(HttpUriParse parse, HttpUriBuild build) {
CHECK(BuildWith(build, element, 0) == url);
}
// Cyberpunk 2077 (PPSA03974) creates a template with a NULL user agent.
void TestTemplateWithoutUserAgent() {
using namespace Libs::Network;
Initialize();
const int pool = Net::NetPoolCreate("http", 0x10000, 0);
const int ssl = Ssl::SslInit(0x10000);
const int http = Http::HttpInit(pool, ssl, 0x10000);
CHECK(Http::HttpCreateTemplate(http, nullptr, 2, 1) > 0);
Shutdown();
}
} // namespace
int main() {
@@ -176,5 +188,6 @@ int main() {
TestAbsentQuery(parse);
TestEmptyUri(parse);
TestPresentQuery(parse);
TestTemplateWithoutUserAgent();
return failures == 0 ? 0 : 1;
}
+342
View File
@@ -0,0 +1,342 @@
#include "libs/libJpegDec.cpp"
#define STB_IMAGE_WRITE_IMPLEMENTATION
#define STBI_WRITE_NO_STDIO
#include "stb_image_write.h"
#include <algorithm>
#include <array>
#include <cstdio>
#include <cstdlib>
#include <vector>
namespace Loader {
void SymbolDatabase::Add(const SymbolResolve&, uint64_t, const std::string&) {}
namespace Timer {
double GetTimeMs() {
return 0.0;
}
} // namespace Timer
} // namespace Loader
namespace {
using namespace Libs::JpegDec;
#define CHECK(condition) \
do { \
if (!(condition)) { \
std::fprintf(stderr, "JpegDecTests:%d: %s\n", __LINE__, #condition); \
std::abort(); \
} \
} while (false)
constexpr uint32_t WIDTH = 16;
constexpr uint32_t HEIGHT = 8;
// cjpeg -grayscale -progressive -optimize -quality 100 of an 8x8 image with pixel x*16 + y*8,
// without the JFIF APP0 segment.
constexpr uint8_t PROGRESSIVE[] = {
0xff, 0xd8, 0xff, 0xdb, 0x00, 0x43, 0x00, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01,
0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0x01, 0xff, 0xc2, 0x00, 0x0b, 0x08, 0x00, 0x08, 0x00, 0x08,
0x01, 0x01, 0x11, 0x00, 0xff, 0xc4, 0x00, 0x14, 0x00, 0x01, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x08, 0xff, 0xda, 0x00, 0x08, 0x01, 0x01,
0x00, 0x00, 0x00, 0x01, 0x27, 0xff, 0x00, 0xff, 0xc4, 0x00, 0x16, 0x10, 0x00, 0x03, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x06, 0x07, 0xff,
0xda, 0x00, 0x08, 0x01, 0x01, 0x00, 0x01, 0x05, 0x02, 0x9b, 0x2d, 0x9f, 0xff, 0xc4, 0x00, 0x18,
0x10, 0x00, 0x02, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x03, 0x22, 0x41, 0x52, 0xff, 0xda, 0x00, 0x08, 0x01, 0x01, 0x00, 0x06, 0x3f, 0x02,
0x44, 0x33, 0x47, 0xff, 0xc4, 0x00, 0x16, 0x10, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x51, 0xf0, 0xff, 0xda, 0x00, 0x08, 0x01,
0x01, 0x00, 0x01, 0x3f, 0x21, 0x9c, 0x47, 0xff, 0xda, 0x00, 0x08, 0x01, 0x01, 0x00, 0x00, 0x00,
0x10, 0x7f, 0xff, 0xc4, 0x00, 0x16, 0x10, 0x00, 0x03, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00,
0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0x00, 0xb1, 0xf0, 0xff, 0xda, 0x00, 0x08, 0x01, 0x01,
0x00, 0x01, 0x3f, 0x10, 0x8f, 0x43, 0xff, 0xd9};
// Red 128, green by row, blue by column.
int Expected(uint32_t x, uint32_t y, uint32_t c) {
const int rgb[3] = {128, static_cast<int>(y * 32), static_cast<int>(x * 16)};
return rgb[c];
}
int Luma(uint32_t x, uint32_t y) {
return (77 * 128 + 150 * Expected(x, y, 1) + 29 * Expected(x, y, 2) + 128) >> 8;
}
void CheckNear(int value, int expected) {
CHECK(std::abs(value - expected) <= 8);
}
// Quality 100 keeps 4:4:4 sampling; 50 subsamples to 4:2:0.
std::vector<uint8_t> Encode(int quality, uint32_t width = WIDTH) {
std::vector<uint8_t> rgb(width * HEIGHT * 3);
for (uint32_t i = 0; i < rgb.size(); i++) {
rgb[i] = static_cast<uint8_t>(Expected(i / 3 % width, i / 3 / width, i % 3));
}
std::vector<uint8_t> jpeg;
stbi_write_jpg_to_func(
[](void* context, void* data, int size) {
auto* out = static_cast<std::vector<uint8_t>*>(context);
const auto* bytes = static_cast<const uint8_t*>(data);
out->insert(out->end(), bytes, bytes + size);
},
&jpeg, static_cast<int>(width), HEIGHT, 3, rgb.data(), quality);
return jpeg;
}
// The first marker of this type in an stb_image_write output.
auto Marker(std::vector<uint8_t>& jpeg, uint8_t type) {
const uint8_t marker[2] = {0xff, type};
const auto it = std::search(jpeg.begin(), jpeg.end(), std::begin(marker), std::end(marker));
CHECK(it != jpeg.end());
return it;
}
struct Decoder {
JpegDecContext context {};
void* handle = nullptr;
explicit Decoder(uint32_t max_width) {
const JpegDecCreateParam create {sizeof(JpegDecCreateParam),
JPEG_DEC_ATTRIBUTE_CSC_STD_SAMPLING, max_width};
CHECK(JpegDecQueryMemorySize(&create) == sizeof(context));
CHECK(JpegDecCreate(&create, &context, sizeof(context), &handle) == 0);
}
~Decoder() {
CHECK(JpegDecDelete(handle) == 0);
CHECK(JpegDecDelete(handle) == JPEG_DEC_ERROR_INVALID_HANDLE);
}
};
void TestParseHeader() {
auto jpeg = Encode(100);
JpegDecParseParam parse {jpeg.data(), static_cast<uint32_t>(jpeg.size()), 0, 2};
JpegDecImageInfo info {};
CHECK(JpegDecParseHeader(&parse, &info) == 0);
CHECK(info.image_width == WIDTH && info.image_height == HEIGHT);
CHECK(info.color_space == JPEG_DEC_COLOR_SPACE_YCC && info.component_number == 3);
CHECK(info.component_id[0] == 1 && info.component_id[1] == 2 && info.component_id[2] == 3);
CHECK(info.sampling_factor[0] == 0x11 && info.sampling_factor[2] == 0x11);
CHECK(info.coefficient_mem_size == 0);
CHECK(info.suitable_csc_attribute == JPEG_DEC_ATTRIBUTE_CSC_STD_SAMPLING);
CHECK(info.output_image_width == WIDTH / 2 && info.output_image_height == HEIGHT / 2);
auto subsampled = Encode(50);
parse = {subsampled.data(), static_cast<uint32_t>(subsampled.size()), 1, 8};
CHECK(JpegDecParseHeader(&parse, &info) == 0);
CHECK(info.sampling_factor[0] == 0x22 && info.sampling_factor[1] == 0x11);
CHECK(info.output_image_width == 2 && info.output_image_height == 1);
// 4x1 luma needs full sampling.
Marker(subsampled, 0xc0)[11] = 0x41;
CHECK(JpegDecParseHeader(&parse, &info) == 0);
CHECK(info.suitable_csc_attribute == JPEG_DEC_ATTRIBUTE_CSC_FULL_SAMPLING);
// A progressive frame keeps 16-bit coefficients for 2x1 blocks of each of the 3 components.
const auto sof = Marker(jpeg, 0xc0);
sof[1] = 0xc2;
parse = {jpeg.data(), static_cast<uint32_t>(jpeg.size()), 0, 1};
CHECK(JpegDecParseHeader(&parse, &info) == 0 && info.coefficient_mem_size == 3 * 2 * 128);
sof[1] = 0xc3;
CHECK(JpegDecParseHeader(&parse, &info) == JPEG_DEC_ERROR_UNSUPPORT_DATA);
parse.down_scale = 3;
CHECK(JpegDecParseHeader(&parse, &info) == JPEG_DEC_ERROR_INVALID_PARAM);
parse = {jpeg.data(), 0, 0, 1};
CHECK(JpegDecParseHeader(&parse, &info) == JPEG_DEC_ERROR_INVALID_SIZE);
CHECK(JpegDecParseHeader(&parse, nullptr) == JPEG_DEC_ERROR_INVALID_ADDR);
jpeg[1] = 0;
parse = {jpeg.data(), static_cast<uint32_t>(jpeg.size()), 0, 1};
CHECK(JpegDecParseHeader(&parse, &info) == JPEG_DEC_ERROR_INVALID_DATA);
}
void TestDecode() {
const auto jpeg = Encode(100);
Decoder decoder(WIDTH);
alignas(4) std::array<uint8_t, (WIDTH * 4 + 8) * HEIGHT> image {};
// GT7 passes a non-NULL coefficient address with size 0 for baseline images.
JpegDecDecodeParam param {.jpeg_mem_addr = jpeg.data(),
.image_mem_addr = image.data(),
.coefficient_mem_addr = image.data(),
.jpeg_mem_size = static_cast<uint32_t>(jpeg.size()),
.image_mem_size = image.size(),
.down_scale = 1,
.pixel_format = JPEG_DEC_PIXEL_FORMAT_R8G8B8A8,
.alpha_value = 200,
.image_pitch = WIDTH * 4 + 8};
image.fill(0xee);
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) ==
static_cast<int32_t>(WIDTH << 16 | HEIGHT));
for (uint32_t y = 0; y < HEIGHT; y++) {
const auto* row = &image[y * param.image_pitch];
for (uint32_t x = 0; x < WIDTH; x++) {
for (uint32_t c = 0; c < 3; c++) {
CheckNear(row[x * 4 + c], Expected(x, y, c));
}
CHECK(row[x * 4 + 3] == 200);
}
CHECK(row[WIDTH * 4] == 0xee && row[WIDTH * 4 + 7] == 0xee);
}
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_B8G8R8A8;
param.image_pitch = 0;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
CheckNear(image[(3 * WIDTH + 5) * 4], Expected(5, 3, 2));
CheckNear(image[(3 * WIDTH + 5) * 4 + 2], 128);
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_Y8;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
CheckNear(image[3 * WIDTH + 5], Luma(5, 3));
// YUYV holds both lumas and the pair's average chroma; saturated pairs stay at 255.
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_Y8U8Y8V8;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
const auto* pair = &image[(3 * WIDTH + 4) * 2];
CheckNear(pair[0], Luma(4, 3));
CheckNear(pair[2], Luma(5, 3));
const int blue = (Expected(4, 3, 2) + Expected(5, 3, 2)) / 2;
CheckNear(pair[1], (-43 * 128 - 85 * Expected(4, 3, 1) + 128 * blue + 32896) >> 8);
CheckNear(pair[3], (128 * 128 - 107 * Expected(4, 3, 1) - 21 * blue + 32896) >> 8);
constexpr uint8_t red[3] = {255, 0, 0};
constexpr uint8_t pure_blue[3] = {0, 0, 255};
CHECK(Chroma(pure_blue, pure_blue, -43, -85, 128) == 255);
CHECK(Chroma(red, red, 128, -107, -21) == 255);
// Each downscaled pixel averages its 4x4 block: (1, 1) covers columns and rows 4-7.
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_R8G8B8A8;
param.down_scale = 4;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) ==
static_cast<int32_t>(WIDTH / 4 << 16 | HEIGHT / 4));
CheckNear(image[(WIDTH / 4 + 1) * 4 + 1], (4 + 5 + 6 + 7) * 32 / 4);
CheckNear(image[(WIDTH / 4 + 1) * 4 + 2], (4 + 5 + 6 + 7) * 16 / 4);
// An odd YUYV width repeats the last pixel.
const auto odd = Encode(100, WIDTH - 1);
param.jpeg_mem_addr = odd.data();
param.jpeg_mem_size = static_cast<uint32_t>(odd.size());
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_Y8U8Y8V8;
param.down_scale = 1;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
const auto* last = &image[(3 * WIDTH + WIDTH - 2) * 2];
CheckNear(last[0], Luma(WIDTH - 2, 3));
CHECK(last[2] == last[0]);
}
void TestProgressive() {
Decoder decoder(8);
std::array<uint8_t, 128> coefficients {};
alignas(4) std::array<uint8_t, 8 * 8 * 4> image {};
JpegDecDecodeParam param {.jpeg_mem_addr = PROGRESSIVE,
.image_mem_addr = image.data(),
.coefficient_mem_addr = coefficients.data(),
.jpeg_mem_size = sizeof(PROGRESSIVE),
.image_mem_size = image.size(),
.coefficient_mem_size = coefficients.size(),
.down_scale = 1,
.pixel_format = JPEG_DEC_PIXEL_FORMAT_R8G8B8A8,
.alpha_value = 255};
JpegDecImageInfo info {};
CHECK(JpegDecDecode(decoder.handle, &param, &info) == (8 << 16 | 8));
CHECK(info.color_space == JPEG_DEC_COLOR_SPACE_GRAYSCALE && info.component_number == 1);
for (uint32_t i = 0; i < 64; i++) {
CheckNear(image[i * 4], i % 8 * 16 + i / 8 * 8);
}
param.coefficient_mem_size--;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_COEF_MEMORY);
param.coefficient_mem_addr = nullptr;
param.coefficient_mem_size = coefficients.size();
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_COEF_MEMORY);
}
// A scan whose Huffman tables no DHT defined gets zeroed tables, not the previous decode's: they
// consume no bits, so stb fails at the scan's first stuffed 0xff.
void TestMissingTables() {
auto jpeg = Encode(100);
Decoder decoder(WIDTH);
alignas(4) std::array<uint8_t, WIDTH * HEIGHT * 4> image {};
JpegDecDecodeParam param {.jpeg_mem_addr = jpeg.data(),
.image_mem_addr = image.data(),
.jpeg_mem_size = static_cast<uint32_t>(jpeg.size()),
.image_mem_size = image.size(),
.down_scale = 1,
.pixel_format = JPEG_DEC_PIXEL_FORMAT_R8G8B8A8,
.alpha_value = 255};
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
const auto dht = Marker(jpeg, 0xc4);
jpeg.erase(dht, dht + 2 + (dht[2] << 8 | dht[3]));
param.jpeg_mem_size = static_cast<uint32_t>(jpeg.size());
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_DECODE_ERROR);
}
void TestDecodeValidation() {
auto jpeg = Encode(100);
Decoder decoder(WIDTH);
alignas(4) std::array<uint8_t, WIDTH * HEIGHT * 4> image {};
const JpegDecDecodeParam valid {.jpeg_mem_addr = jpeg.data(),
.image_mem_addr = image.data(),
.jpeg_mem_size = static_cast<uint32_t>(jpeg.size()),
.image_mem_size = image.size(),
.down_scale = 1,
.pixel_format = JPEG_DEC_PIXEL_FORMAT_R8G8B8A8,
.alpha_value = 255};
CHECK(JpegDecDecode(decoder.handle, &valid, nullptr) > 0);
CHECK(JpegDecDecode(nullptr, &valid, nullptr) == JPEG_DEC_ERROR_INVALID_HANDLE);
CHECK(JpegDecDecode(decoder.handle, nullptr, nullptr) == JPEG_DEC_ERROR_INVALID_ADDR);
auto param = valid;
param.image_mem_addr = image.data() + 2;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_ADDR);
param.pixel_format = JPEG_DEC_PIXEL_FORMAT_Y8;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) > 0);
param = valid;
param.image_mem_size--;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_SIZE);
param = valid;
param.image_pitch = WIDTH * 4 - 4;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_PARAM);
param = valid;
param.pixel_format = 2;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_PARAM);
param = valid;
param.alpha_value = 256;
CHECK(JpegDecDecode(decoder.handle, &param, nullptr) == JPEG_DEC_ERROR_INVALID_PARAM);
// The width limit applies after downscaling.
Decoder narrow(WIDTH / 2);
CHECK(JpegDecDecode(narrow.handle, &valid, nullptr) == JPEG_DEC_ERROR_INVALID_WORK_MEMORY);
param = valid;
param.down_scale = 2;
CHECK(JpegDecDecode(narrow.handle, &param, nullptr) > 0);
// Full sampling needs a decoder created for it.
const auto sof = Marker(jpeg, 0xc0);
sof[11] = 0x41;
CHECK(JpegDecDecode(decoder.handle, &valid, nullptr) == JPEG_DEC_ERROR_INVALID_WORK_MEMORY);
// An output above 32767 pixels needs image information.
sof[11] = 0x11;
sof[7] = 40000 >> 8;
sof[8] = 40000 & 0xff;
Decoder wide(0xffff);
CHECK(JpegDecDecode(wide.handle, &valid, nullptr) == JPEG_DEC_ERROR_INVALID_ADDR);
JpegDecCreateParam create {sizeof(JpegDecCreateParam), 3, WIDTH};
CHECK(JpegDecQueryMemorySize(&create) == JPEG_DEC_ERROR_INVALID_PARAM);
create = {sizeof(JpegDecCreateParam), JPEG_DEC_ATTRIBUTE_CSC_FULL_SAMPLING, 0};
CHECK(JpegDecQueryMemorySize(&create) == JPEG_DEC_ERROR_INVALID_PARAM);
create.this_size = 8;
CHECK(JpegDecQueryMemorySize(&create) == JPEG_DEC_ERROR_INVALID_SIZE);
CHECK(JpegDecQueryMemorySize(nullptr) == JPEG_DEC_ERROR_INVALID_ADDR);
}
} // namespace
int main() {
TestParseHeader();
TestDecode();
TestProgressive();
TestMissingTables();
TestDecodeValidation();
std::puts("JPEG decoder tests passed");
}
+205
View File
@@ -0,0 +1,205 @@
#include "libs/libJpegEnc.cpp"
#define STB_IMAGE_IMPLEMENTATION
#define STBI_ONLY_JPEG
#define STBI_NO_STDIO
#include "stb_image.h"
#include <array>
#include <cstdio>
#include <cstdlib>
namespace Loader {
void SymbolDatabase::Add(const SymbolResolve&, uint64_t, const std::string&) {}
namespace Timer {
double GetTimeMs() {
return 0.0;
}
} // namespace Timer
} // namespace Loader
namespace {
using namespace Libs::JpegEnc;
#define CHECK(condition) \
do { \
if (!(condition)) { \
std::fprintf(stderr, "JpegEncTests:%d: %s\n", __LINE__, #condition); \
std::abort(); \
} \
} while (false)
constexpr uint32_t WIDTH = 16;
constexpr uint32_t HEIGHT = 8;
using Jpeg = std::array<uint8_t, 4096>;
using Pattern = int (*)(uint32_t x, uint32_t y, uint32_t c);
struct Encoder {
JpegEncContext context {};
void* handle = nullptr;
Encoder() {
const JpegEncCreateParam create {sizeof(JpegEncCreateParam), 0};
CHECK(JpegEncQueryMemorySize(&create) == sizeof(context));
CHECK(JpegEncCreate(&create, &context, sizeof(context), &handle) == 0);
}
~Encoder() {
CHECK(JpegEncDelete(handle) == 0);
CHECK(JpegEncDelete(handle) == JPEG_ENC_ERROR_INVALID_HANDLE);
}
};
// Red 128, green by row, blue by column.
int Expected(uint32_t x, uint32_t y, uint32_t c) {
const int rgb[3] = {128, static_cast<int>(y * 32), static_cast<int>(x * 16)};
return rgb[c];
}
int Luma(uint32_t x, uint32_t y, uint32_t /*c*/) {
return (77 * 128 + 150 * Expected(x, y, 1) + 29 * Expected(x, y, 2) + 128) >> 8;
}
// Decodes the output and checks every channel against the pattern.
void CheckDecoded(const Jpeg& jpeg, int32_t size, Pattern pattern) {
CHECK(size > 0);
int width = 0;
int height = 0;
auto* rgb = stbi_load_from_memory(jpeg.data(), size, &width, &height, nullptr, 3);
CHECK(rgb != nullptr && width == WIDTH && height == HEIGHT);
for (uint32_t i = 0; i < WIDTH * HEIGHT * 3; i++) {
CHECK(std::abs(rgb[i] - pattern(i / 3 % WIDTH, i / 3 / WIDTH, i % 3)) <= 8);
}
stbi_image_free(rgb);
}
void TestEncode() {
Encoder encoder;
alignas(4) std::array<uint8_t, WIDTH * HEIGHT * 4> bgra {};
std::array<uint8_t, WIDTH * HEIGHT> y8 {};
for (uint32_t i = 0; i < WIDTH * HEIGHT; i++) {
const uint32_t x = i % WIDTH;
const uint32_t y = i / WIDTH;
bgra[i * 4] = Expected(x, y, 2);
bgra[i * 4 + 1] = Expected(x, y, 1);
bgra[i * 4 + 2] = Expected(x, y, 0);
y8[i] = Luma(x, y, 0);
}
Jpeg jpeg {};
JpegEncEncodeParam param {.image_mem_addr = bgra.data(),
.jpeg_mem_addr = jpeg.data(),
.image_mem_size = bgra.size(),
.jpeg_mem_size = jpeg.size(),
.image_width = WIDTH,
.image_height = HEIGHT,
.image_pitch = WIDTH * 4,
.pixel_format = JPEG_ENC_PIXEL_FORMAT_B8G8R8A8,
.color_space = JPEG_ENC_COLOR_SPACE_YCC,
.sampling_type = JPEG_ENC_SAMPLING_TYPE_420};
JpegEncOutputInfo info {};
const auto size = JpegEncEncode(encoder.handle, &param, &info);
CHECK(info.data_size == static_cast<uint32_t>(size) && info.processed_height == HEIGHT);
CheckDecoded(jpeg, size, Expected);
const auto color = param;
// The same pixels read as RGBA swap red and blue.
param.pixel_format = JPEG_ENC_PIXEL_FORMAT_R8G8B8A8;
CheckDecoded(jpeg, JpegEncEncode(encoder.handle, &param, nullptr),
[](uint32_t x, uint32_t y, uint32_t c) { return Expected(x, y, 2 - c); });
param = color;
param.color_space = JPEG_ENC_COLOR_SPACE_GRAYSCALE;
param.sampling_type = JPEG_ENC_SAMPLING_TYPE_FULL;
CheckDecoded(jpeg, JpegEncEncode(encoder.handle, &param, nullptr), Luma);
param.image_mem_addr = y8.data();
param.image_mem_size = y8.size();
param.image_pitch = WIDTH;
param.pixel_format = JPEG_ENC_PIXEL_FORMAT_Y8;
CheckDecoded(jpeg, JpegEncEncode(encoder.handle, &param, nullptr), Luma);
// GT7's YUYV rows hold 2 bytes per pixel at a pitch of width * 4.
constexpr uint8_t pair[4] = {64, 160, 192, 96};
alignas(4) std::array<uint8_t, WIDTH * HEIGHT * 4> yuyv {};
for (uint32_t i = 0; i < yuyv.size(); i++) {
yuyv[i] = i % (WIDTH * 4) < WIDTH * 2 ? pair[i % 4] : 0xff;
}
param = color;
param.image_mem_addr = yuyv.data();
param.pixel_format = JPEG_ENC_PIXEL_FORMAT_Y8U8Y8V8;
param.sampling_type = JPEG_ENC_SAMPLING_TYPE_422;
CheckDecoded(jpeg, JpegEncEncode(encoder.handle, &param, nullptr),
[](uint32_t x, uint32_t, uint32_t c) {
constexpr int rgb[2][3] = {{19, 76, 121}, {147, 204, 249}};
return rgb[x % 2][c];
});
param.color_space = JPEG_ENC_COLOR_SPACE_GRAYSCALE;
param.sampling_type = JPEG_ENC_SAMPLING_TYPE_FULL;
CheckDecoded(jpeg, JpegEncEncode(encoder.handle, &param, nullptr),
[](uint32_t x, uint32_t, uint32_t) { return x % 2 == 0 ? 64 : 192; });
param = color;
param.jpeg_mem_size = 100;
CHECK(JpegEncEncode(encoder.handle, &param, &info) == JPEG_ENC_ERROR_DATA_OVERFLOW);
CHECK(info.data_size <= 100 && info.processed_height == 0);
}
void TestValidation() {
JpegEncCreateParam create {sizeof(JpegEncCreateParam), 1};
CHECK(JpegEncQueryMemorySize(nullptr) == JPEG_ENC_ERROR_INVALID_ADDR);
CHECK(JpegEncQueryMemorySize(&create) == JPEG_ENC_ERROR_INVALID_PARAM);
create = {4, 0};
CHECK(JpegEncQueryMemorySize(&create) == JPEG_ENC_ERROR_INVALID_SIZE);
Encoder encoder;
alignas(4) std::array<uint8_t, WIDTH * HEIGHT * 4> rgba {};
Jpeg jpeg {};
const JpegEncEncodeParam valid {.image_mem_addr = rgba.data(),
.jpeg_mem_addr = jpeg.data(),
.image_mem_size = rgba.size(),
.jpeg_mem_size = jpeg.size(),
.image_width = WIDTH,
.image_height = HEIGHT,
.image_pitch = WIDTH * 4,
.pixel_format = JPEG_ENC_PIXEL_FORMAT_R8G8B8A8,
.color_space = JPEG_ENC_COLOR_SPACE_YCC,
.sampling_type = JPEG_ENC_SAMPLING_TYPE_422,
.compression_ratio = 50,
.restart_interval = -1};
CHECK(JpegEncEncode(encoder.handle, &valid, nullptr) > 0);
CHECK(JpegEncEncode(nullptr, &valid, nullptr) == JPEG_ENC_ERROR_INVALID_HANDLE);
CHECK(JpegEncEncode(encoder.handle, nullptr, nullptr) == JPEG_ENC_ERROR_INVALID_ADDR);
auto param = valid;
param.image_mem_addr = rgba.data() + 1;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_ADDR);
param = valid;
param.image_mem_size--;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_SIZE);
param = valid;
param.image_pitch = WIDTH * 4 - 4;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
param = valid;
param.sampling_type = JPEG_ENC_SAMPLING_TYPE_FULL;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
param = valid;
param.pixel_format = JPEG_ENC_PIXEL_FORMAT_Y8;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
param = valid;
param.restart_interval = -2;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
param = valid;
param.encode_mode = 2;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
param = valid;
param.pixel_format = 2;
CHECK(JpegEncEncode(encoder.handle, &param, nullptr) == JPEG_ENC_ERROR_INVALID_PARAM);
}
} // namespace
int main() {
TestEncode();
TestValidation();
std::puts("JPEG encoder tests passed");
}
+122
View File
@@ -0,0 +1,122 @@
#include "gameViews.h"
#include "launcherTheme.h"
#include "ui_configuration_list_widget.h"
#include <QApplication>
#include <QColor>
#include <QDir>
#include <QEvent>
#include <QImage>
#include <QScrollBar>
#include <QTemporaryDir>
#include <cstdio>
#include <cstdlib>
namespace {
void Check(bool value, const char *message) {
if (!value) {
std::fprintf(stderr, "LauncherGameViewsTests: failed: %s\n", message);
std::abort();
}
}
// The red pic0 dominates a pixel only where nothing opaque is painted over the background.
bool ShowsBackground(const QColor &color) {
return color.red() > color.green() + 40 && color.red() > color.blue() + 40;
}
// The card panel is the app's (white) Window color over the background.
bool ShowsPanel(const QColor &color) {
return !ShowsBackground(color) &&
qAbs(color.lightness() - QApplication::palette().color(QPalette::Window).lightness()) < 32;
}
class PaintCounter final : public QObject {
public:
int paints = 0;
protected:
bool eventFilter(QObject * /*watched*/, QEvent *event) override {
paints += static_cast<int>(event->type() == QEvent::Paint);
return false;
}
};
} // namespace
int main(int argc, char **argv) {
qputenv("QT_QPA_PLATFORM", "offscreen");
QApplication application(argc, argv);
// Breeze, where installed, paints its own color behind scroll bars; Fusion's are opaque.
const bool breeze = QApplication::setStyle(QStringLiteral("Breeze")) != nullptr;
// A fixed light palette, whatever the desktop colors: a black panel then stands out.
QApplication::setPalette(QPalette(Qt::white));
LauncherTheme::Initialize(application);
QTemporaryDir game;
Check(QDir(game.path()).mkpath(QStringLiteral("sce_sys")), "cannot create sce_sys");
QImage pic0(64, 36, QImage::Format_RGB32);
pic0.fill(Qt::red);
Check(pic0.save(game.filePath(QStringLiteral("sce_sys/pic0.png"))), "cannot write pic0.png");
QWidget host;
Ui::ConfigurationListWidget ui;
ui.setupUi(&host);
auto *views = ui.game_views;
auto *list = ui.cfgs_list;
auto *grid = ui.cfgs_grid;
SetupGameGrid(list, grid);
grid->setIconSize(QSize(64, 64));
QPixmap cover(64, 64);
cover.fill(Qt::blue);
auto *item = new QTreeWidgetItem(list, {QStringLiteral("Game")});
item->setIcon(GAME_NAME_COLUMN, QIcon(cover));
host.resize(480, 360);
views->SetBackgroundImage(game.path());
host.show();
QApplication::processEvents();
const QPoint empty(views->width() - 60, views->height() - 30);
Check(ShowsBackground(views->grab().toImage().pixelColor(empty)),
"the game background is hidden behind the list view");
views->setCurrentWidget(grid);
QApplication::processEvents();
const auto image = views->grab().toImage();
Check(ShowsBackground(image.pixelColor(empty)),
"the game background is hidden behind the grid view");
const auto card = grid->visualRect(grid->model()->index(0, GAME_NAME_COLUMN))
.translated(grid->viewport()->mapTo(views, QPoint()));
Check(card.width() == 64 + 16, "the card is not the cover plus its padding wide");
Check(image.pixelColor(card.topLeft() + QPoint(8 + 32, 8 + 32)) == QColor(Qt::blue),
"the cover is not drawn inside the card padding");
Check(ShowsPanel(image.pixelColor(card.topLeft() + QPoint(4, 40))),
"the card panel does not back its padding");
PaintCounter counter;
grid->viewport()->installEventFilter(&counter);
item->setText(GAME_SIZE_COLUMN, QStringLiteral("1.00 GB"));
QApplication::processEvents();
Check(counter.paints > 0, "the card does not repaint when the game size arrives");
if (breeze) {
for (QAbstractScrollArea *view : {static_cast<QAbstractScrollArea *>(grid),
static_cast<QAbstractScrollArea *>(list)}) {
view->setVerticalScrollBarPolicy(Qt::ScrollBarAlwaysOn);
views->setCurrentWidget(view);
QApplication::processEvents();
const auto *bar = view->verticalScrollBar();
Check(ShowsBackground(views->grab().toImage().pixelColor(
bar->mapTo(views, QPoint(2, bar->height() - 20)))),
"Breeze hides the game background behind a scroll bar");
}
}
std::printf("LauncherGameViewsTests: passed%s\n", breeze ? "" : " (Breeze checks skipped)");
return 0;
}
+55
View File
@@ -0,0 +1,55 @@
#include "graphics/host_gpu/renderer/occlusionQueries.h"
#include <array>
#include <cstdio>
#include <cstdlib>
namespace {
using namespace Libs::Graphics;
#define CHECK(condition) \
do { \
if (!(condition)) { \
std::fprintf(stderr, "OcclusionQueriesTests:%d: %s\n", __LINE__, #condition); \
std::abort(); \
} \
} while (false)
using Results = std::array<uint64_t, OCCLUSION_QUERY_DBS * 2>;
// The samples a query passed, as ZPass predication sums them; -1 before every counter is written.
int64_t Passed(const Results& results) {
int64_t samples = 0;
for (uint32_t db = 0; db < OCCLUSION_QUERY_DBS; db++) {
const auto begin = results[db * 2];
const auto end = results[db * 2 + 1];
if ((begin & end & OCCLUSION_QUERY_READY) == 0) {
return -1;
}
samples += static_cast<int64_t>(end - begin);
}
return samples;
}
void TestCounters() {
Results results {};
WriteZPassCounters(results.data(), 100, true);
CHECK(Passed(results) == -1);
WriteZPassCounters(results.data() + 1, 100, false);
CHECK(Passed(results) == 0);
WriteZPassCounters(results.data() + 1, 142, false);
CHECK(Passed(results) == 42);
CHECK(results[0] == (OCCLUSION_QUERY_READY | 100) && results[30] == OCCLUSION_QUERY_READY);
// A begin clears the end a previous query wrote late, so the query is not ready until its own.
WriteZPassCounters(results.data(), 142, true);
CHECK(Passed(results) == -1);
WriteZPassCounters(results.data() + 1, 150, false);
CHECK(Passed(results) == 8);
}
} // namespace
int main() {
TestCounters();
std::puts("Occlusion query tests passed");
}
+124
View File
@@ -436,6 +436,128 @@ void TestInvariantIndirectImageMaterialization(bool nested) {
"malformed image table tracking was not transactional");
}
void TestResidentImageHeap() {
struct FakeImageHeap final : ImageHeapSource {
bool accept = true;
std::vector<DescriptorValue> kinds;
std::vector<ImageHeapRequest> requests;
bool Acquire(const ImageHeapRequest &request, ImageHeapView &view) override {
requests.push_back(request);
if (!accept) return false;
view = {.slots = 0x123456789abcull, .slot_count = 36u, .kinds = kinds};
return true;
}
};
auto fixture = MakeIndirectImageFixture(false, 48u, 0u, 16u);
fixture->PlanAndTrack();
const auto plan = ExtractResourcePlan(fixture->program);
std::array<uint32_t, 8> descriptor{};
descriptor[0] = 0x20u;
descriptor[1] = static_cast<uint32_t>(Libs::Graphics::Prospero::BufferFormat::k32_32_32_32Float) << 20u;
descriptor[2] = 3u | (3u << 14u);
descriptor[3] = Libs::Graphics::DstSel(4, 5, 6, 7) |
(static_cast<uint32_t>(Libs::Graphics::Prospero::ImageType::kColor2D) << 28u);
auto cube = descriptor;
cube[3] = Libs::Graphics::DstSel(4, 5, 6, 7) |
(static_cast<uint32_t>(Libs::Graphics::Prospero::ImageType::kCube) << 28u);
std::array<uint32_t, 16> data{0x3000u, 4u << 16u, 0x100000u, 0u,
0x2000u, 48u << 16u, 3u, 0u};
std::copy_n(data.begin(), 4u, data.begin() + 12u);
LinearTestMemory memory;
memory.fail_address = 0x3000u;
memory.watched_address = 0x2000u;
std::copy(descriptor.begin(), descriptor.end(), memory.words.begin() + 0x1010u / 4u);
FakeImageHeap heap;
heap.kinds = {{descriptor, 8u}, {cube, 8u}};
SrtRuntime runtime{.user_data = data, .userdata = &memory,
.read_specialization_memory = ReadLinearTestMemory, .image_heap = &heap};
ResourceSnapshot snapshot;
ResourceSpecialization specialization;
Check(MaterializeResources(plan, runtime, snapshot, specialization) &&
heap.requests.size() == 1u && heap.requests[0].base == 0x2000u &&
heap.requests[0].size == 144u && heap.requests[0].step == 16u &&
heap.requests[0].first == 0u && heap.requests[0].count == 9u &&
heap.requests[0].root == &plan.info.images[0] &&
memory.watched_reads == 0u && snapshot.images.size() == 3u,
"resident image heap was materialized entry by entry");
const auto &root = specialization.images[0];
const auto mapping = root.indirect_mapping_offset;
Check(root.heap && root.indirect_root == 0u && root.indirect_search_iterations == 0u &&
snapshot.flattened_srt.size() == mapping + 3u &&
snapshot.flattened_srt[mapping] == 36u &&
snapshot.flattened_srt[mapping + 1u] == 0x56789abcu &&
snapshot.flattened_srt[mapping + 2u] == 0x1234u &&
specialization.images[1].heap && specialization.images[1].indirect_root == 0u &&
!specialization.images[1].cube && specialization.images[2].heap &&
specialization.images[2].cube && snapshot.images[2].dwords == cube &&
std::ranges::find(snapshot.specialization_reads,
std::pair<uint64_t, uint64_t>{0x2000u, 144u}) !=
snapshot.specialization_reads.end(),
"resident image heap lost its slot mapping, typed kinds or read dependency");
auto &program = fixture->program;
ApplyResourceSpecialization(program, specialization);
ShaderComputeInputInfo compute{};
CollectShaderInfo(program, {.compute = &compute});
Check(program.info.uses_image_heap && program.info.images.size() == 3u &&
program.info.images[0].indirect_resources == std::vector<uint32_t>({0u, 1u, 2u}) &&
std::ranges::all_of(program.info.images, &ImageResource::heap),
"applied heap kinds lost their root or heap flag");
heap.kinds.clear();
Check(MaterializeResources(plan, runtime, snapshot, specialization) &&
snapshot.images.size() == 1u && !specialization.images[0].heap &&
specialization.images[0].indirect_root == ImageResource::NoIndirectImage,
"a resident heap without kinds did not read as one null image");
heap.accept = false;
Check(MaterializeResources(plan, runtime, snapshot, specialization) &&
heap.requests.size() == 3u && memory.watched_reads == 1u &&
snapshot.images.size() == 2u && !specialization.images[0].heap &&
snapshot.images[1].dwords == descriptor,
"a declined heap was not materialized entry by entry");
}
void TestImageHeapKinds() {
using Format = Libs::Graphics::Prospero::BufferFormat;
using Type = Libs::Graphics::Prospero::ImageType;
ImageResource root;
root.resource_class = ImageResourceClass::Sampled;
root.dimension = Decoder::ImageDimension::Dim2DArray;
root.source = 3u;
const auto descriptor = [](Format format, Type type) {
DescriptorValue value{.dword_count = 8u};
value.dwords[0] = 0x20u;
value.dwords[1] = static_cast<uint32_t>(format) << 20u;
value.dwords[2] = 3u | (3u << 14u);
value.dwords[3] = Libs::Graphics::DstSel(4, 5, 6, 7) | (static_cast<uint32_t>(type) << 28u);
return value;
};
const auto kind_of = [&](Format format, Type type) {
auto value = descriptor(format, type);
return ImageHeapKind(root, value);
};
const auto float_kind = kind_of(Format::k32_32_32_32Float, Type::kColor2DArray);
const auto uint_kind = kind_of(Format::k32UInt, Type::kColor2DArray);
const auto cube_kind = kind_of(Format::k32_32_32_32Float, Type::kCube);
auto volume = descriptor(Format::k32_32_32_32Float, Type::kColor3D);
Check(float_kind != 0u && uint_kind != 0u && cube_kind != 0u && float_kind != uint_kind &&
float_kind != cube_kind && kind_of(Format::kInvalid, Type::kColor2DArray) == 0u &&
ImageHeapKind(root, volume) == 0u &&
std::ranges::all_of(volume.dwords, [](uint32_t word) { return word == 0u; }),
"heap kinds did not separate typed entries or null mismatched ones");
const auto resource = ImageHeapResource(root, uint_kind);
const auto cube_resource = ImageHeapResource(root, cube_kind);
auto other = root;
other.source = 9u;
Check(resource.numeric_class == Libs::Graphics::Prospero::TextureNumericClass::Uint &&
!resource.cube && cube_resource.cube &&
cube_resource.numeric_class == Libs::Graphics::Prospero::TextureNumericClass::Float &&
resource.resource_class == ImageResourceClass::Sampled && resource.read &&
ImageHeapResource(root, 0u) == ImageHeapResource(other, 0u),
"heap kind resources lost their type or kept root identity");
root.depth_compare = true;
Check(kind_of(Format::k32UInt, Type::kColor2DArray) == 0u,
"depth comparison accepted an integer heap entry");
}
void TestBoundedImageViewEligibility() {
using Type = Libs::Graphics::Prospero::ImageType;
auto fixture = MakeIndirectImageFixture(false, 48u, 0u, 16u);
@@ -8796,6 +8918,8 @@ int main() {
Run("dynamic storage mips", TestDynamicStorageMipTracking);
Run("invariant indirect images", [] { TestInvariantIndirectImageMaterialization(true); });
Run("unbounded material image keys", [] { TestInvariantIndirectImageMaterialization(false); });
Run("resident image heap", TestResidentImageHeap);
Run("image heap kinds", TestImageHeapKinds);
Run("bounded image view eligibility", TestBoundedImageViewEligibility);
Run("exact unfiltered image candidates", TestExactUnfilteredImageCandidates);
Run("waterfall image table", TestWaterfallImageTable);
File diff suppressed because it is too large. Load diff
+38
View File
@@ -2049,6 +2049,36 @@ void TestDirectReleaseContracts() {
free_offset, SceKernelPageSize) == Libs::LibKernel::KERNEL_ERROR_ENOENT,
"checked release did not report an unallocated range");
// An unchecked release frees the allocated parts on both sides of a gap.
const auto page = static_cast<int64_t>(SceKernelPageSize);
for (const auto offset: {free_offset - page * 3, free_offset - page}) {
int64_t allocated = 0;
CheckOk(test,
Libs::LibKernel::Memory::KernelAllocateDirectMemory(
offset, offset + page, SceKernelPageSize, SceKernelPageSize, SceKernelMtypeC,
&allocated),
"KernelAllocateDirectMemory(around a gap)");
Check(test, allocated == offset, "allocation did not land at its fixed offset");
}
for (const auto [offset, len]:
{std::pair {free_offset - page * 3, page * 3}, {free_offset - page, page * 2}}) {
void* mapping = nullptr;
Check(test,
Libs::LibKernel::Memory::KernelMapDirectMemory(&mapping, len, SceKernelProtCpuRw, 0,
offset, SceKernelPageSize) ==
Libs::LibKernel::KERNEL_ERROR_ENOMEM,
"direct map accepted a gap or the unallocated tail");
}
CheckOk(test,
Libs::LibKernel::Memory::KernelReleaseDirectMemory(free_offset - page * 3, page * 3),
"KernelReleaseDirectMemory(over a gap)");
for (const auto offset: {free_offset - page * 3, free_offset - page}) {
Check(test,
Libs::LibKernel::Memory::KernelCheckedReleaseDirectMemory(offset, page) ==
Libs::LibKernel::KERNEL_ERROR_ENOENT,
"unchecked release left an allocation around a gap");
}
std::printf("[host] %-48s ok\n", test);
}
@@ -3127,6 +3157,14 @@ void TestAutomaticMemoryFragmentedMapRollback() {
"automatic allocation while second donation remains aliased");
CheckOk(test, KernelMunmap(reinterpret_cast<uint64_t>(alias), block * 2),
"unmap mixed ownership alias");
CheckOk(
test,
KernelMapDirectMemory(&alias, block * 2, SceKernelProtCpuRw, SceKernelMapFixed, gap, block),
"alias over a free donation behind ordinary ownership");
CheckFailed(test, MapAutomaticMemory(base, block * 2, SceKernelMtypeC, SceKernelProtCpuRw),
"automatic allocation of a donation claimed by a new alias");
CheckOk(test, KernelMunmap(reinterpret_cast<uint64_t>(alias), block * 2),
"unmap claiming alias");
CheckFailed(test, MapAutomaticMemory(base, block * 3, SceKernelMtypeC, SceKernelProtCpuRw),
"automatic allocation beyond donated capacity");
CheckOk(test, MapAutomaticMemory(base, block * 2, SceKernelMtypeC, SceKernelProtCpuRw),
+16
View File
@@ -7045,6 +7045,17 @@ void TestPixelAncillaryLayerInput() {
CheckSpirvBinaryValidates(wqm_result.spirv);
}
uint32_t dpp[] = {
EncodeVop3Word0(0x148, 6), EncodeVop3Word1(5 + 256, 128 + 16, 128 + 11),
EncodeVop1(0x01, 5, 250), EncodeVop1Dpp(6) | (1u << 19u), // Bound-zero DPP overwrite.
EncodeExp0(0x00, 0x1), EncodeExp1(5, 0, 0, 0), EncodeSopp(0x01),
};
auto dpp_result = RecompileForTest(dpp, options);
Check(ProgramHasInput(dpp_result.program, StageInputKind::Layer) &&
!ProgramHasInput(dpp_result.program, StageInputKind::PackedAncillary),
"DPP write retained the old ancillary register value");
CheckSpirvBinaryValidates(dpp_result.spirv);
const uint32_t loop[] = {
EncodeSop1(0x04, 8, 126), // Save the original EXEC mask.
EncodeSop1(0x0a, 126, 8), // Widen for helper lanes.
@@ -7098,6 +7109,9 @@ void TestPixelAncillaryLayerInput() {
shader[5] = EncodeExp1(6, 7, 5, 0);
ExpectFatal([&] { (void)RecompileForTest(shader, options); },
"direct raw ancillary export was silently replaced after field lowering");
dpp[3] = EncodeVop1Dpp(6); // Without BOUND_CTRL, lanes with inactive sources keep the value.
ExpectFatal([&] { (void)RecompileForTest(dpp, options); },
"DPP write without BOUND_CTRL silently replaced the ancillary value");
#endif
}
@@ -11146,6 +11160,8 @@ void TestMeshExportStorage() {
!SpirvSourceHasInstructionUsing(source, "OpAccessChain",
"vsharp %uint_0 %uint_4294967295"),
"mesh draw prefix was lost or spilled shader data wrapped into push constants");
Check(SpirvSourceHasInstructionUsing(source, "OpBitFieldUExtract", "%uint_10 %uint_9"),
"NGG primitive vertex indices must be 9 bits; bit 9 of each field is the edge flag");
const auto &binary = result.spirv;
for (const auto builtin : {3u, 4u}) {
auto expected = builtin == 3u ? std::vector<uint32_t>{0, 3}