mirror/PSPRecomp
1
0
Fork 0
mirror of https://github.com/jessicanataliagta/PSPRecomp synced 2026-09-26 08:41:08 -04:00
Code Issues Packages Projects Releases Wiki Activity

otimizações round 7

Browse Source 
otimizações round 7
This commit is contained in:
Jessica_Natalia
2026-08-16 23:12:34 -03:00
parent 839a6f2c32
commit b70c9b6bfb
15 changed files with 3137 additions and 2538 deletions
Download Patch File Download Diff File Expand all files Collapse all files
profiles/vcs/CMakeLists.txt
+9 -2
View File
@@ -151,7 +151,7 @@ if(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang")
set_source_files_properties("${VCS_PROFILE_DIR}/host/vcs_profile.cpp" PROPERTIES COMPILE_OPTIONS "-O2")
endif()
# Tier-2 V2 keeps each profile-guided cluster in its own translation unit.
# Tier-2 V4 keeps each profile-guided cluster in its own translation unit.
# This is the actual second AOT layer: hot functions/loops are cold-split from
# the giant generated units and selected cross-unit edges are fused locally.
# /GL- on these sources avoids re-inflating them into a giant LTCG partition;
@@ -170,11 +170,18 @@ if(MSVC)
"${VCS_PROFILE_DIR}/host/vcs_tier2_superblocks.cpp"
${VCS_TIER2_CLUSTER_SOURCES}
PROPERTIES COMPILE_OPTIONS "/O2;/Ob3;/bigobj;/GL-")
# Geometry is the largest V4 hot TU (~7k generated lines) and /Ob3 can
# make MSVC spend many minutes in pathological inlining/optimization.
# Keep full /O2 codegen and normal aggressive /Ob2 inlining here; this
# preserves GPR shadow/SIMD/dataflow while avoiding the compile-time cliff.
set_source_files_properties(
"${VCS_PROFILE_DIR}/host/vcs_tier2_cluster_geometry.cpp"
PROPERTIES COMPILE_OPTIONS "/O2;/Ob2;/bigobj;/GL-")
elseif(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang")
set_source_files_properties(
"${VCS_PROFILE_DIR}/host/vcs_tier2_superblocks.cpp"
${VCS_TIER2_CLUSTER_SOURCES}
PROPERTIES COMPILE_OPTIONS "-O3;-g0")
PROPERTIES COMPILE_OPTIONS "-O2;-g0")
endif()
set(VCS_GPU_BACKEND_SOURCE host/ge_gpu_backend_dx12.cpp)
profiles/vcs/host/ge_gpu_backend.hpp
+7
View File
@@ -374,6 +374,13 @@ struct GeGpuBackendReport {
std::uint64_t dx12_batch_appends{};
std::uint64_t dx12_batch_merges{};
std::uint64_t dx12_gpu_draw_calls{};
// Tier-2 V4 DX12: adjacent state-compatible draws can be submitted through
// ExecuteIndirect. The GPU still executes every PSP draw, but CPU command
// recording collapses transform/pixel constant updates + Draw* calls into
// one API command for the whole run.
std::uint64_t dx12_indirect_executes{};
std::uint64_t dx12_indirect_draws{};
std::uint64_t dx12_indirect_saved_api_draws{};
std::uint64_t dx12_srv_high_water{};
std::uint64_t dx12_gpu_feedback_draws{};
std::uint64_t dx12_self_feedback_snapshots{};
profiles/vcs/host/ge_gpu_backend_dx12.cpp
+305 -70
View File
@@ -44,6 +44,10 @@ using Microsoft::WRL::ComPtr;
constexpr std::uint32_t kReferenceWidth = 480u;
constexpr std::uint32_t kReferenceHeight = 272u;
constexpr std::size_t kGeometryUploadCapacity = 64u * 1024u * 1024u;
// V4 ExecuteIndirect arguments live in their own persistently mapped upload
// arena. Keeping them separate means a draw-heavy frame can never steal bytes
// from the established 64 MiB geometry budget. 4 MiB holds >20k commands.
constexpr std::size_t kIndirectUploadCapacity = 4u * 1024u * 1024u;
// Stage 45.2: persistent per-frame texture upload arena. The 44.7 path created,
// mapped and destroyed one committed upload resource for every decoded texture.
// Streaming bursts therefore paid kernel/D3D12 allocation overhead on the hot GE
@@ -122,6 +126,23 @@ struct Dx12PixelConstants {
};
static_assert(sizeof(Dx12PixelConstants) == 5u * sizeof(std::uint32_t));
// Tier-2 V4 / 150-FPS path. ExecuteIndirect moves the two per-draw root
// constant writes and Draw* call out of the CPU command-recording loop. Runs
// still preserve guest order; only adjacent draws with identical fixed GPU
// state participate.
struct Dx12IndirectDrawCommand {
std::uint32_t transform[40]{};
std::uint32_t pixel[5]{};
D3D12_DRAW_ARGUMENTS draw{};
};
struct Dx12IndirectDrawIndexedCommand {
std::uint32_t transform[40]{};
std::uint32_t pixel[5]{};
D3D12_DRAW_INDEXED_ARGUMENTS draw{};
};
static_assert(sizeof(Dx12IndirectDrawCommand) == 196u);
static_assert(sizeof(Dx12IndirectDrawIndexedCommand) == 200u);
struct CloudCameraCandidate {
std::array<float, 12> view{};
std::array<float, 16> projection{};
@@ -162,6 +183,8 @@ struct Dx12FrameResources {
ComPtr<ID3D12CommandAllocator> allocator;
ComPtr<ID3D12Resource> upload_buffer;
std::byte *mapped_upload{};
ComPtr<ID3D12Resource> indirect_upload_buffer;
std::byte *mapped_indirect_upload{};
ComPtr<ID3D12Resource> texture_upload_buffer;
std::byte *mapped_texture_upload{};
std::size_t texture_upload_cursor{};
@@ -274,6 +297,8 @@ struct Dx12GeState {
UINT64 next_fence{1u};
ComPtr<ID3D12RootSignature> root_signature;
ComPtr<ID3D12RootSignature> cloud_root_signature;
ComPtr<ID3D12CommandSignature> indirect_draw_signature;
ComPtr<ID3D12CommandSignature> indirect_draw_indexed_signature;
ComPtr<ID3DBlob> vertex_shader;
ComPtr<ID3DBlob> packed_0115_vertex_shader;
ComPtr<ID3DBlob> pixel_shader;
@@ -388,6 +413,10 @@ std::uint64_t texture_key(const GeGpuDrawDescriptor &draw) noexcept {
return key;
}
// Forward declaration used by batch/indirect compatibility checks. The sampler
// implementation lives next to sampler creation below.
std::uint64_t sampler_key(const GeGpuDrawDescriptor &draw) noexcept;
D3D12_CPU_DESCRIPTOR_HANDLE rtv_cpu(Dx12GeState &s, UINT index) noexcept {
D3D12_CPU_DESCRIPTOR_HANDLE h = s.rtv_heap->GetCPUDescriptorHandleForHeapStart();
h.ptr += static_cast<SIZE_T>(index) * s.rtv_stride;
@@ -1043,7 +1072,9 @@ bool adjacent_batch_merge_compatible(const Dx12Batch &a, const Dx12Batch &b) noe
(b.hardware_transform && b.transform.primitive != 3u)) return false;
if (pipeline_key(a.draw) != pipeline_key(b.draw)) return false;
if (a.draw.texture_enabled != b.draw.texture_enabled) return false;
if (a.draw.texture_enabled && texture_key(a.draw) != texture_key(b.draw)) return false;
if (a.draw.texture_enabled &&
(texture_key(a.draw) != texture_key(b.draw) || sampler_key(a.draw) != sampler_key(b.draw)))
return false;
if (a.draw.scissor_x0 != b.draw.scissor_x0 || a.draw.scissor_y0 != b.draw.scissor_y0 ||
a.draw.scissor_x1 != b.draw.scissor_x1 || a.draw.scissor_y1 != b.draw.scissor_y1)
return false;
@@ -1060,6 +1091,107 @@ bool adjacent_batch_merge_compatible(const Dx12Batch &a, const Dx12Batch &b) noe
return !a.hardware_transform || hardware_transform_equal(a.transform, b.transform);
}
bool dx12_execute_indirect_enabled() noexcept {
static const bool enabled = [] {
const char *text = std::getenv("PSPRECOMP_DX12_EXECUTE_INDIRECT");
return text == nullptr || (*text != '\0' && std::strcmp(text, "0") != 0 &&
std::strcmp(text, "false") != 0 && std::strcmp(text, "FALSE") != 0 &&
std::strcmp(text, "off") != 0 && std::strcmp(text, "OFF") != 0);
}();
return enabled;
}
D3D12_PRIMITIVE_TOPOLOGY batch_topology(const Dx12Batch &batch) noexcept {
return batch.hardware_transform && batch.transform.primitive == 4u
? D3D_PRIMITIVE_TOPOLOGY_TRIANGLESTRIP
: D3D_PRIMITIVE_TOPOLOGY_TRIANGLELIST;
}
std::uint64_t batch_full_pipeline_key(const Dx12Batch &batch) noexcept {
const bool cull = batch.hardware_transform && batch.transform.cull_enabled;
const bool accept_ccw = cull && batch.transform.accept_counter_clockwise;
return pipeline_key(batch.draw) |
(batch.packed_0115 ? (std::uint64_t{1} << 63u) : 0u) |
(cull ? (std::uint64_t{1} << 62u) : 0u) |
(accept_ccw ? (std::uint64_t{1} << 61u) : 0u);
}
bool indirect_run_compatible(const Dx12Batch &a, const Dx12Batch &b) noexcept {
if (a.framebuffer_feedback || b.framebuffer_feedback) return false;
if (a.draw.clear_mode || b.draw.clear_mode) return false;
if (a.indexed != b.indexed || a.packed_0115 != b.packed_0115) return false;
if ((a.draw.framebuffer_address & 0x001FFFF0u) !=
(b.draw.framebuffer_address & 0x001FFFF0u)) return false;
if (batch_topology(a) != batch_topology(b)) return false;
if (batch_full_pipeline_key(a) != batch_full_pipeline_key(b)) return false;
if (a.draw.texture_enabled != b.draw.texture_enabled) return false;
if (a.draw.texture_enabled &&
(texture_key(a.draw) != texture_key(b.draw) || sampler_key(a.draw) != sampler_key(b.draw)))
return false;
if (a.draw.scissor_x0 != b.draw.scissor_x0 || a.draw.scissor_y0 != b.draw.scissor_y0 ||
a.draw.scissor_x1 != b.draw.scissor_x1 || a.draw.scissor_y1 != b.draw.scissor_y1)
return false;
const Dx12BlendPlan ba = dx12_blend_plan(a.draw);
const Dx12BlendPlan bb = dx12_blend_plan(b.draw);
if (ba.uses_constant != bb.uses_constant) return false;
if (ba.uses_constant && ba.constant_rgb != bb.constant_rgb) return false;
return true;
}
void account_executed_batch(Dx12GeState &s, const Dx12Batch &batch,
std::uint32_t srv_index, const Dx12BlendPlan &blend_plan) {
if (batch.draw.depth_test_enabled) s.report.depth_tested_game_draw_calls += batch.logical_draw_count;
if (batch.draw.depth_write_enabled) s.report.depth_writing_game_draw_calls += batch.logical_draw_count;
if (batch.draw.alpha_test_enabled) s.report.alpha_tested_game_draw_calls += batch.logical_draw_count;
switch (blend_variant(batch.draw)) {
case 1u: s.report.standard_alpha_blended_game_draw_calls += batch.logical_draw_count; break;
case 2u: s.report.fixed_replace_blended_game_draw_calls += batch.logical_draw_count; break;
case 3u: s.report.additive_blended_game_draw_calls += batch.logical_draw_count; break;
default: break;
}
if (batch.draw.blend_enabled && !batch.draw.clear_mode && !blend_plan.exact) {
s.report.unsupported_blend_game_draw_calls += batch.logical_draw_count;
static std::uint32_t diagnostic_count = 0u;
if (std::getenv("PSPRECOMP_DX12_BLEND_DIAG") != nullptr && diagnostic_count < 32u) {
std::ostringstream line;
line << "DX12 unsupported PSP blend fallback #" << (diagnostic_count + 1u)
<< ": eq=" << (batch.draw.blend_equation & 7u)
<< " src=" << (batch.draw.blend_source_factor & 0xFu)
<< " dst=" << (batch.draw.blend_dest_factor & 0xFu)
<< " fixS=0x" << std::hex << (batch.draw.blend_fix_source & 0x00FFFFFFu)
<< " fixD=0x" << (batch.draw.blend_fix_dest & 0x00FFFFFFu) << std::dec;
const std::string message = line.str();
std::cerr << "[blend] " << message << "\n";
runtime_log_error("blend", message);
++diagnostic_count;
}
}
if (batch.draw.fog_enabled) s.report.fogged_game_draw_calls += batch.logical_draw_count;
if (srv_index != 0u) {
const std::uint32_t submitted_vertices = batch.indexed ? batch.index_count : batch.vertex_count;
s.report.textured_game_triangles +=
batch.hardware_transform && batch.transform.primitive == 4u
? (submitted_vertices > 2u ? submitted_vertices - 2u : 0u)
: submitted_vertices / 3u;
s.report.textured_game_vertices += submitted_vertices;
switch (batch.draw.texture_function & 7u) {
case 0u: s.report.modulate_texture_game_draw_calls += batch.logical_draw_count; break;
case 1u: s.report.decal_texture_game_draw_calls += batch.logical_draw_count; break;
case 2u: s.report.blend_texture_game_draw_calls += batch.logical_draw_count; break;
case 3u: s.report.replace_texture_game_draw_calls += batch.logical_draw_count; break;
case 4u: s.report.add_texture_game_draw_calls += batch.logical_draw_count; break;
default: ++s.report.unsupported_texture_function_game_draw_calls; break;
}
if (batch.draw.texture_double_color)
s.report.double_color_texture_game_draw_calls += batch.logical_draw_count;
if (batch.draw.texture_mipmap_enabled) {
s.report.mipmapped_game_draw_calls += batch.logical_draw_count;
if (batch.draw.texture_mipmap_linear)
s.report.mip_linear_game_draw_calls += batch.logical_draw_count;
}
}
}
bool append_or_merge_batch(Dx12GeState &s, Dx12Batch batch) {
static const bool merge_enabled = [] {
const char *text = std::getenv("PSPRECOMP_DX12_BATCH_MERGE");
@@ -1456,6 +1588,42 @@ bool create_root_signature(Dx12GeState &s, std::string &error) noexcept {
return true;
}
bool create_indirect_signatures(Dx12GeState &s, std::string &error) noexcept {
std::array<D3D12_INDIRECT_ARGUMENT_DESC, 3> args{};
args[0].Type = D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT;
args[0].Constant.RootParameterIndex = 2u;
args[0].Constant.DestOffsetIn32BitValues = 0u;
args[0].Constant.Num32BitValuesToSet = 40u;
args[1].Type = D3D12_INDIRECT_ARGUMENT_TYPE_CONSTANT;
args[1].Constant.RootParameterIndex = 3u;
args[1].Constant.DestOffsetIn32BitValues = 0u;
args[1].Constant.Num32BitValuesToSet = 5u;
D3D12_COMMAND_SIGNATURE_DESC desc{};
desc.NumArgumentDescs = static_cast<UINT>(args.size());
desc.pArgumentDescs = args.data();
args[2].Type = D3D12_INDIRECT_ARGUMENT_TYPE_DRAW;
desc.ByteStride = sizeof(Dx12IndirectDrawCommand);
HRESULT hr = s.device->CreateCommandSignature(&desc, s.root_signature.Get(),
IID_PPV_ARGS(&s.indirect_draw_signature));
if (FAILED(hr)) {
error = hr_text(hr, "CreateCommandSignature(DX12 GE draw)");
return false;
}
args[2].Type = D3D12_INDIRECT_ARGUMENT_TYPE_DRAW_INDEXED;
desc.ByteStride = sizeof(Dx12IndirectDrawIndexedCommand);
hr = s.device->CreateCommandSignature(&desc, s.root_signature.Get(),
IID_PPV_ARGS(&s.indirect_draw_indexed_signature));
if (FAILED(hr)) {
error = hr_text(hr, "CreateCommandSignature(DX12 GE draw indexed)");
s.indirect_draw_signature.Reset();
return false;
}
return true;
}
bool create_targets(Dx12GeState &s, std::string &error) noexcept {
D3D12_DESCRIPTOR_HEAP_DESC rtv_desc{};
rtv_desc.Type = D3D12_DESCRIPTOR_HEAP_TYPE_RTV;
@@ -2283,7 +2451,8 @@ const CloudCameraCandidate *select_cloud_camera(const Dx12GeState &s) noexcept {
bool changed = true;
while (changed && ancestors.size() < kFramebufferTargetCapacity) {
changed = false;
for (const Dx12Batch &batch : s.batches) {
for (std::size_t batch_cursor = 0u; batch_cursor < s.batches.size(); ++batch_cursor) {
const Dx12Batch &batch = s.batches[batch_cursor];
if (!batch.framebuffer_feedback) continue;
const std::uint32_t source = batch.feedback_address & 0x001FFFF0u;
const std::uint32_t destination = batch.draw.framebuffer_address & 0x001FFFF0u;
@@ -3080,6 +3249,25 @@ bool create_backend(Dx12GeState &s, std::string &error) noexcept {
if (FAILED(hr) || mapped == nullptr) { error = hr_text(hr, "Map(DX12 GE geometry upload)"); return false; }
frame.mapped_upload = static_cast<std::byte *>(mapped);
// Optional V4 arena. Failure here must never make the renderer fail to
// boot on an older/low-memory driver; that frame simply uses scalar
// Draw*/root-constant recording.
D3D12_RESOURCE_DESC indirect_upload = upload;
indirect_upload.Width = kIndirectUploadCapacity;
hr = s.device->CreateCommittedResource(&upload_heap, D3D12_HEAP_FLAG_NONE, &indirect_upload,
D3D12_RESOURCE_STATE_GENERIC_READ, nullptr,
IID_PPV_ARGS(&frame.indirect_upload_buffer));
if (SUCCEEDED(hr) && frame.indirect_upload_buffer) {
mapped = nullptr;
hr = frame.indirect_upload_buffer->Map(0u, &no_read, &mapped);
if (SUCCEEDED(hr) && mapped != nullptr) {
frame.mapped_indirect_upload = static_cast<std::byte *>(mapped);
} else {
frame.indirect_upload_buffer.Reset();
frame.mapped_indirect_upload = nullptr;
}
}
if (s.texture_upload_ring_enabled) {
D3D12_RESOURCE_DESC texture_upload = upload;
texture_upload.Width = kTextureUploadCapacity;
@@ -3113,6 +3301,16 @@ bool create_backend(Dx12GeState &s, std::string &error) noexcept {
if (s.fence_event == nullptr) { error = "CreateEventW failed for DX12 GE fence"; return false; }
if (!compile_shaders(s, error)) return false;
if (!create_root_signature(s, error)) return false;
{
std::string indirect_error;
if (!create_indirect_signatures(s, indirect_error)) {
// Scalar recording is the fully supported fallback. Command
// signatures are an optimization, never a backend requirement.
s.indirect_draw_signature.Reset();
s.indirect_draw_indexed_signature.Reset();
if (!indirect_error.empty()) runtime_log_error("dx12 execute indirect disabled", indirect_error);
}
}
if (!create_cloud_root_signature(s, error)) return false;
if (!create_targets(s, error)) return false;
if (!create_present_pipeline(s, error)) return false;
@@ -3144,11 +3342,15 @@ void destroy_backend(Dx12GeState &s) noexcept {
for (Dx12FrameResources &frame : s.frames) {
if (frame.upload_buffer && frame.mapped_upload != nullptr)
frame.upload_buffer->Unmap(0u, nullptr);
if (frame.indirect_upload_buffer && frame.mapped_indirect_upload != nullptr)
frame.indirect_upload_buffer->Unmap(0u, nullptr);
if (frame.texture_upload_buffer && frame.mapped_texture_upload != nullptr)
frame.texture_upload_buffer->Unmap(0u, nullptr);
frame.mapped_upload = nullptr;
frame.mapped_indirect_upload = nullptr;
frame.mapped_texture_upload = nullptr;
frame.transient_resources.clear();
frame.indirect_upload_buffer.Reset();
frame.texture_upload_buffer.Reset();
frame.upload_buffer.Reset();
frame.texture_upload_cursor = 0u;
@@ -3186,6 +3388,8 @@ void destroy_backend(Dx12GeState &s) noexcept {
s.pixel_shader.Reset();
s.packed_0115_vertex_shader.Reset();
s.vertex_shader.Reset();
s.indirect_draw_signature.Reset();
s.indirect_draw_indexed_signature.Reset();
s.root_signature.Reset();
s.cloud_root_signature.Reset();
s.readback_buffer.Reset();
@@ -3904,6 +4108,9 @@ bool ge_gpu_backend_finish_color_frame(std::uint64_t vblank) noexcept {
Dx12FramebufferTarget *display_target = find_framebuffer_target(s, s.display_framebuffer);
Dx12FrameResources &frame = s.frames[s.frame_cursor];
const bool indirect_enabled = dx12_execute_indirect_enabled() &&
s.indirect_draw_signature && s.indirect_draw_indexed_signature &&
frame.indirect_upload_buffer && frame.mapped_indirect_upload != nullptr;
std::string error;
if (!wait_for_fence(s, frame.fence_value, error)) {
runtime_log_error("dx12 ge frame wait", error);
@@ -3974,6 +4181,7 @@ bool ge_gpu_backend_finish_color_frame(std::uint64_t vblank) noexcept {
Dx12FramebufferTarget *current_target = nullptr;
std::uint32_t current_address = 0xFFFFFFFFu;
std::uint32_t executed_batches = 0u;
std::size_t indirect_cursor = 0u;
std::uint32_t bound_srv = std::numeric_limits<std::uint32_t>::max();
std::uint32_t bound_sampler = std::numeric_limits<std::uint32_t>::max();
Dx12TransformConstants active_transform{};
@@ -4006,7 +4214,8 @@ bool ge_gpu_backend_finish_color_frame(std::uint64_t vblank) noexcept {
}
constexpr float black[4]{0.0f, 0.0f, 0.0f, 1.0f};
for (const Dx12Batch &batch : s.batches) {
for (std::size_t batch_cursor = 0u; batch_cursor < s.batches.size(); ++batch_cursor) {
const Dx12Batch &batch = s.batches[batch_cursor];
const std::uint32_t address = batch.draw.framebuffer_address & 0x001FFFF0u;
const bool cloud_target_batch = address == cloud_target_address;
if (trace_cloud_frame && (cloud_target_batch || address == s.display_framebuffer)) {
@@ -4192,26 +4401,45 @@ bool ge_gpu_backend_finish_color_frame(std::uint64_t vblank) noexcept {
bound_sampler = sampler_index;
}
// Discover an indirect run before writing root constants. If the run
// fits the dedicated arena, ExecuteIndirect owns those constants too,
// avoiding even the first scalar SetGraphicsRoot32BitConstants pair.
std::size_t indirect_end = batch_cursor + 1u;
if (indirect_enabled && !trace_cloud_frame && !batch.framebuffer_feedback &&
!batch.draw.clear_mode) {
while (indirect_end < s.batches.size() &&
indirect_run_compatible(batch, s.batches[indirect_end])) {
++indirect_end;
}
}
const std::size_t indirect_count = indirect_end - batch_cursor;
const std::size_t indirect_stride = batch.indexed
? sizeof(Dx12IndirectDrawIndexedCommand)
: sizeof(Dx12IndirectDrawCommand);
const std::size_t indirect_command_bytes = indirect_stride * indirect_count;
const bool execute_indirect_run = indirect_count >= 3u &&
indirect_cursor + indirect_command_bytes <= kIndirectUploadCapacity;
const std::uint32_t logical_width = current_target != nullptr && current_target->logical_width != 0u
? current_target->logical_width : kReferenceWidth;
const std::uint32_t logical_height = current_target != nullptr && current_target->logical_height != 0u
? current_target->logical_height : kReferenceHeight;
const Dx12TransformConstants draw_transform =
make_transform_constants(batch, logical_width, logical_height);
if (!active_transform_valid ||
std::memcmp(&draw_transform, &active_transform, sizeof(draw_transform)) != 0) {
s.list->SetGraphicsRoot32BitConstants(
2u, 40u, &draw_transform, 0u);
active_transform = draw_transform;
active_transform_valid = true;
}
const Dx12PixelConstants pixel_state = make_pixel_constants(batch.draw, srv_index != 0u);
if (!active_pixel_valid ||
std::memcmp(&pixel_state, &active_pixel, sizeof(pixel_state)) != 0) {
s.list->SetGraphicsRoot32BitConstants(3u, 5u, &pixel_state, 0u);
active_pixel = pixel_state;
active_pixel_valid = true;
if (!execute_indirect_run) {
const Dx12TransformConstants draw_transform =
make_transform_constants(batch, logical_width, logical_height);
const Dx12PixelConstants pixel_state = make_pixel_constants(batch.draw, srv_index != 0u);
if (!active_transform_valid ||
std::memcmp(&draw_transform, &active_transform, sizeof(draw_transform)) != 0) {
s.list->SetGraphicsRoot32BitConstants(2u, 40u, &draw_transform, 0u);
active_transform = draw_transform;
active_transform_valid = true;
}
if (!active_pixel_valid ||
std::memcmp(&pixel_state, &active_pixel, sizeof(pixel_state)) != 0) {
s.list->SetGraphicsRoot32BitConstants(3u, 5u, &pixel_state, 0u);
active_pixel = pixel_state;
active_pixel_valid = true;
}
}
// The PSP scissor is expressed in 480x272 logical pixels and has to be
@@ -4266,63 +4494,70 @@ bool ge_gpu_backend_finish_color_frame(std::uint64_t vblank) noexcept {
active_blend_fix = fix;
}
}
// V4 high-margin path: collapse adjacent draws that differ only in
// per-draw root constants into one ExecuteIndirect call. No draw is
// reordered and framebuffer-feedback/clear boundaries never participate.
if (execute_indirect_run) {
const std::size_t stride = indirect_stride;
const std::size_t command_bytes = indirect_command_bytes;
std::byte *command_dst = frame.mapped_indirect_upload + indirect_cursor;
for (std::size_t j = batch_cursor; j < indirect_end; ++j) {
const Dx12Batch &ibatch = s.batches[j];
const Dx12TransformConstants itransform =
make_transform_constants(ibatch, logical_width, logical_height);
const Dx12PixelConstants ipixel =
make_pixel_constants(ibatch.draw, srv_index != 0u);
if (batch.indexed) {
Dx12IndirectDrawIndexedCommand command{};
std::memcpy(command.transform, &itransform, sizeof(itransform));
std::memcpy(command.pixel, &ipixel, sizeof(ipixel));
command.draw.IndexCountPerInstance = ibatch.index_count;
command.draw.InstanceCount = 1u;
command.draw.StartIndexLocation = ibatch.first_index;
command.draw.BaseVertexLocation = static_cast<INT>(ibatch.first_vertex);
command.draw.StartInstanceLocation = 0u;
std::memcpy(command_dst, &command, sizeof(command));
} else {
Dx12IndirectDrawCommand command{};
std::memcpy(command.transform, &itransform, sizeof(itransform));
std::memcpy(command.pixel, &ipixel, sizeof(ipixel));
command.draw.VertexCountPerInstance = ibatch.vertex_count;
command.draw.InstanceCount = 1u;
command.draw.StartVertexLocation = ibatch.first_vertex;
command.draw.StartInstanceLocation = 0u;
std::memcpy(command_dst, &command, sizeof(command));
}
command_dst += stride;
}
ID3D12CommandSignature *signature = batch.indexed
? s.indirect_draw_indexed_signature.Get()
: s.indirect_draw_signature.Get();
s.list->ExecuteIndirect(signature, static_cast<UINT>(indirect_count),
frame.indirect_upload_buffer.Get(), indirect_cursor,
nullptr, 0u);
indirect_cursor += command_bytes;
executed_batches += static_cast<std::uint32_t>(indirect_count);
++s.report.dx12_indirect_executes;
s.report.dx12_indirect_draws += indirect_count;
s.report.dx12_indirect_saved_api_draws += indirect_count - 1u;
for (std::size_t j = batch_cursor; j < indirect_end; ++j)
account_executed_batch(s, s.batches[j], srv_index, blend_plan);
batch_cursor = indirect_end - 1u;
cloud_batch_index += indirect_count - 1u;
// ExecuteIndirect leaves root constants equal to the last command;
// force the scalar cache to repopulate before the next ordinary draw.
active_transform_valid = false;
active_pixel_valid = false;
continue;
}
if (batch.indexed)
s.list->DrawIndexedInstanced(batch.index_count, 1u, batch.first_index,
static_cast<INT>(batch.first_vertex), 0u);
else
s.list->DrawInstanced(batch.vertex_count, 1u, batch.first_vertex, 0u);
++executed_batches;
if (batch.draw.depth_test_enabled) s.report.depth_tested_game_draw_calls += batch.logical_draw_count;
if (batch.draw.depth_write_enabled) s.report.depth_writing_game_draw_calls += batch.logical_draw_count;
if (batch.draw.alpha_test_enabled) s.report.alpha_tested_game_draw_calls += batch.logical_draw_count;
switch (blend_variant(batch.draw)) {
case 1u: s.report.standard_alpha_blended_game_draw_calls += batch.logical_draw_count; break;
case 2u: s.report.fixed_replace_blended_game_draw_calls += batch.logical_draw_count; break;
case 3u: s.report.additive_blended_game_draw_calls += batch.logical_draw_count; break;
default: break;
}
if (batch.draw.blend_enabled && !batch.draw.clear_mode && !blend_plan.exact) {
s.report.unsupported_blend_game_draw_calls += batch.logical_draw_count;
static std::uint32_t diagnostic_count = 0u;
if (std::getenv("PSPRECOMP_DX12_BLEND_DIAG") != nullptr && diagnostic_count < 32u) {
std::ostringstream line;
line << "DX12 unsupported PSP blend fallback #" << (diagnostic_count + 1u)
<< ": eq=" << (batch.draw.blend_equation & 7u)
<< " src=" << (batch.draw.blend_source_factor & 0xFu)
<< " dst=" << (batch.draw.blend_dest_factor & 0xFu)
<< " fixS=0x" << std::hex << (batch.draw.blend_fix_source & 0x00FFFFFFu)
<< " fixD=0x" << (batch.draw.blend_fix_dest & 0x00FFFFFFu) << std::dec;
const std::string message = line.str();
std::cerr << "[blend] " << message << "\n";
runtime_log_error("blend", message);
++diagnostic_count;
}
}
if (batch.draw.fog_enabled) s.report.fogged_game_draw_calls += batch.logical_draw_count;
if (srv_index != 0u) {
const std::uint32_t submitted_vertices = batch.indexed ? batch.index_count : batch.vertex_count;
s.report.textured_game_triangles +=
batch.hardware_transform && batch.transform.primitive == 4u
? (submitted_vertices > 2u ? submitted_vertices - 2u : 0u)
: submitted_vertices / 3u;
s.report.textured_game_vertices += submitted_vertices;
switch (batch.draw.texture_function & 7u) {
case 0u: s.report.modulate_texture_game_draw_calls += batch.logical_draw_count; break;
case 1u: s.report.decal_texture_game_draw_calls += batch.logical_draw_count; break;
case 2u: s.report.blend_texture_game_draw_calls += batch.logical_draw_count; break;
case 3u: s.report.replace_texture_game_draw_calls += batch.logical_draw_count; break;
case 4u: s.report.add_texture_game_draw_calls += batch.logical_draw_count; break;
default: ++s.report.unsupported_texture_function_game_draw_calls; break;
}
if (batch.draw.texture_double_color)
s.report.double_color_texture_game_draw_calls += batch.logical_draw_count;
if (batch.draw.texture_mipmap_enabled) {
s.report.mipmapped_game_draw_calls += batch.logical_draw_count;
if (batch.draw.texture_mipmap_linear)
s.report.mip_linear_game_draw_calls += batch.logical_draw_count;
}
}
account_executed_batch(s, batch, srv_index, blend_plan);
}
if (trace_cloud_frame) {
runtime_log_line(std::string("CLOUD_TRACE_END injected=") +
profiles/vcs/host/vcs_fast_paths.hpp
+42
View File
@@ -109,3 +109,45 @@ finish:
}
} // namespace psprecomp
#if defined(__SSE2__) || defined(_M_X64) || defined(_M_IX86_FP)
#include <immintrin.h>
#endif
namespace psprecomp {
// Tier-2 V4: ordered SIMD primitives. Multiplication is vectorized while the
// four products are reduced in the original left-to-right order. This avoids
// FMA/reassociation surprises while removing most scalar multiply issue cost in
// VCS's matrix-heavy geometry hot path.
inline float vcs_tier2_dot4_ordered(const float *a, const float *b) noexcept {
#if defined(__SSE2__) || defined(_M_X64) || (defined(_M_IX86_FP) && _M_IX86_FP >= 2)
const __m128 va = _mm_loadu_ps(a);
const __m128 vb = _mm_loadu_ps(b);
const __m128 vm = _mm_mul_ps(va, vb);
alignas(16) float p[4];
_mm_store_ps(p, vm);
return ((p[0] + p[1]) + p[2]) + p[3];
#else
return ((a[0] * b[0] + a[1] * b[1]) + a[2] * b[2]) + a[3] * b[3];
#endif
}
inline void vcs_tier2_mat4_mul_ordered(const float *s, const float *t, float *d) noexcept {
for (std::uint32_t a = 0; a < 4u; ++a) {
const float *tr = t + a * 4u;
d[a * 4u + 0u] = vcs_tier2_dot4_ordered(s + 0u, tr);
d[a * 4u + 1u] = vcs_tier2_dot4_ordered(s + 4u, tr);
d[a * 4u + 2u] = vcs_tier2_dot4_ordered(s + 8u, tr);
d[a * 4u + 3u] = vcs_tier2_dot4_ordered(s + 12u, tr);
}
}
inline void vcs_tier2_mat4_vec_first3_ordered(const float *m, const float *v,
float *out) noexcept {
out[0] = vcs_tier2_dot4_ordered(m + 0u, v);
out[1] = vcs_tier2_dot4_ordered(m + 4u, v);
out[2] = vcs_tier2_dot4_ordered(m + 8u, v);
}
} // namespace psprecomp
profiles/vcs/host/vcs_profile.cpp
+3
View File
@@ -7600,6 +7600,9 @@ void install_profile(psprecomp::Runtime &runtime, std::uint32_t user_arena_start
<< " gpu_draws=" << d(r.dx12_gpu_draw_calls, o.dx12_gpu_draw_calls)
<< " batch_appends=" << d(r.dx12_batch_appends, o.dx12_batch_appends)
<< " batch_merges=" << d(r.dx12_batch_merges, o.dx12_batch_merges)
<< " indirect_exec=" << d(r.dx12_indirect_executes, o.dx12_indirect_executes)
<< " indirect_draws=" << d(r.dx12_indirect_draws, o.dx12_indirect_draws)
<< " indirect_saved=" << d(r.dx12_indirect_saved_api_draws, o.dx12_indirect_saved_api_draws)
<< " tex_req=" << d(r.texture_decode_requests, o.texture_decode_requests)
<< " tex_hits=" << d(r.texture_cache_hits, o.texture_cache_hits)
<< " tex_uploads=" << d(r.decoded_texture_uploads, o.decoded_texture_uploads)
profiles/vcs/host/vcs_runtime_log.cpp
+9 -3
View File
@@ -3,6 +3,8 @@
#include "vcs_tier2_superblocks.hpp"
#include <chrono>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <fstream>
#include <iomanip>
@@ -63,7 +65,7 @@ void runtime_log_initialize(const VcsConfiguration &configuration) {
return;
}
s.file << "VCSNative runtime log\n";
s.file << "stage=tier2-superblock-v3-dataflow-2026-08-16\n";
s.file << "stage=tier2-v4-150fps-2026-08-16\n";
s.file << "config=" << configuration.source_path.string() << '\n';
s.file << "started=" << timestamp_now() << '\n';
s.file << "perf_telemetry=" << (configuration.diagnostics.perf_telemetry ? 1 : 0)
@@ -72,8 +74,12 @@ void runtime_log_initialize(const VcsConfiguration &configuration) {
s.file << "guest_hotspot=" << (configuration.diagnostics.guest_hotspot_profile ? 1 : 0)
<< " sample_stride=256 interval_vblanks=300\n";
s.file << "tier2_superblocks=" << (tier2_superblocks_enabled() ? 1 : 0)
<< " version=3 clusters=7 mask=0x" << std::hex << tier2_cluster_mask() << std::dec
<< " hot_blocks=1060 static_fused_calls=49 static_fused_tail=13 hooks=25 unwind_fix=1 reentry_guard=1 dataflow=1 vfpu_block32=133 mem_runs=35 mem_words=287 append32=51 advance32=89\n\n";
<< " version=4 clusters=7 mask=0x" << std::hex << tier2_cluster_mask() << std::dec
<< " hot_blocks=1060 static_fused_calls=49 static_fused_tail=13 hooks=25"
<< " unwind_fix=1 reentry_guard=1 dataflow=1 vfpu_block32=133 mem_runs=35 mem_words=287"
<< " append32=51 advance32=89 simd_mat4=4 simd_matvec=19"
<< " gpr_shadow_clusters=4 gpr_shadow_regs=24 gpr_shadow_occurrences=3428 geometry_shadow=0"
<< " dx12_execute_indirect_default=1 indirect_buffer_mb=4\n\n";
if (s.flush_every_line) s.file.flush();
}
profiles/vcs/host/vcs_tier2_cluster_boundary.cpp
+492 -470
View File
File diff suppressed because it is too large Load Diff
profiles/vcs/host/vcs_tier2_cluster_edge43.cpp
+191 -169
View File
@@ -1,8 +1,9 @@
// AUTO-GENERATED by profiles/vcs/tools/build_tier2_superblocks.py.
// Tier-2 SUPERBLOCK V3 DATAFLOW cluster: edge43
// Tier-2 SUPERBLOCK V4 150FPS cluster: edge43
#include "vcs_tier2_superblocks.hpp"
#include "psprecomp/runtime.hpp"
#include "generated_units.hpp"
#include "vcs_fast_paths.hpp"
#include <cstdint>
@@ -27,8 +28,25 @@ void tier2_superblock_edge43(psprecomp::Runtime &rt,
std::uint32_t local_pc = 0u;
std::uint32_t local_transfers = 0u;
std::uint32_t entry_id = 0u;
std::uint32_t tier2_gpr_4 = ctx.gpr[4];
std::uint32_t tier2_gpr_6 = ctx.gpr[6];
std::uint32_t tier2_gpr_5 = ctx.gpr[5];
std::uint32_t tier2_gpr_29 = ctx.gpr[29];
std::uint32_t tier2_gpr_17 = ctx.gpr[17];
std::uint32_t tier2_gpr_19 = ctx.gpr[19];
bool tier2_gpr_shadow_valid = true;
#define TIER2_GPR_SYNC_OUT() do { if (tier2_gpr_shadow_valid) { ctx.gpr[4] = tier2_gpr_4; ctx.gpr[6] = tier2_gpr_6; ctx.gpr[5] = tier2_gpr_5; ctx.gpr[29] = tier2_gpr_29; ctx.gpr[17] = tier2_gpr_17; ctx.gpr[19] = tier2_gpr_19; } } while (false)
#define TIER2_GPR_SYNC_IN() do { if (tier2_gpr_shadow_valid) { tier2_gpr_4 = ctx.gpr[4]; tier2_gpr_6 = ctx.gpr[6]; tier2_gpr_5 = ctx.gpr[5]; tier2_gpr_29 = ctx.gpr[29]; tier2_gpr_17 = ctx.gpr[17]; tier2_gpr_19 = ctx.gpr[19]; } } while (false)
#define TIER2_GPR_BEFORE_COLD() do { TIER2_GPR_SYNC_OUT(); tier2_gpr_shadow_valid = false; } while (false)
#define TIER2_SB_RETURN() do { tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)rt.tier2_complete_fused_transfers(ctx, tier2_tail_count_); --tier2_return_depth; (void)rt.tier2_complete_fused_transfers(ctx, 1u); } if (tier2_pending_transfers != 0u) { (void)rt.tier2_complete_fused_transfers(ctx, tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
auto tier2_complete_shadow = [&](std::uint32_t count) -> bool {
TIER2_GPR_SYNC_OUT();
const bool same = rt.tier2_complete_fused_transfers(ctx, count);
if (!same) tier2_gpr_shadow_valid = false;
return same;
};
#define TIER2_SB_RETURN() do { TIER2_GPR_SYNC_OUT(); tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)tier2_complete_shadow(tier2_tail_count_); --tier2_return_depth; (void)tier2_complete_shadow(1u); } if (tier2_pending_transfers != 0u) { (void)tier2_complete_shadow(tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
goto TIER2_ENTRY_DISPATCH;
@@ -99,11 +117,11 @@ TIER2_LOCAL_DISPATCH_U0043:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -148,11 +166,11 @@ TIER2_LOCAL_DISPATCH_U0044:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -197,8 +215,9 @@ TIER2_ENTRY_DISPATCH:
TIER2_SB_RETURN();
}
// TIER2_GPR_BODY_BEGIN
SB_L_088B1780:
{ const std::uint32_t vfpu_address = ctx.gpr[4] + static_cast<std::uint32_t>(0);
{ const std::uint32_t vfpu_address = tier2_gpr_4 + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -206,7 +225,7 @@ SB_L_088B1780:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<0u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[4] + static_cast<std::uint32_t>(16);
{ const std::uint32_t vfpu_address = tier2_gpr_4 + static_cast<std::uint32_t>(16);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -214,7 +233,7 @@ SB_L_088B1780:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<1u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(0);
{ const std::uint32_t vfpu_address = tier2_gpr_5 + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -222,7 +241,7 @@ SB_L_088B1780:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<4u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(16);
{ const std::uint32_t vfpu_address = tier2_gpr_5 + static_cast<std::uint32_t>(16);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -297,46 +316,46 @@ SB_L_088B17B4:
SB_L_088B180C:
// vflush: architectural no-op that retains VFPU prefixes
ctx.gpr[4] = (ctx.vfpu_scalar_bits_ct<131u>());
ctx.gpr[2] = (ctx.gpr[4] & 16u);
tier2_gpr_4 = (ctx.vfpu_scalar_bits_ct<131u>());
ctx.gpr[2] = (tier2_gpr_4 & 16u);
jump_target = ctx.gpr[31];
ctx.gpr[2] = (ctx.gpr[2] < static_cast<std::uint32_t>(1) ? 1u : 0u);
local_pc = jump_target;
goto TIER2_LOCAL_DISPATCH_U0043;
SB_L_088B18EC:
ctx.gpr[29] = (ctx.gpr[29] + static_cast<std::uint32_t>(-128));
{ const std::uint32_t tier2_words[9]{ctx.gpr[16], ctx.gpr[17], ctx.gpr[18], ctx.gpr[19], ctx.gpr[20], ctx.gpr[21], ctx.gpr[22], ctx.gpr[23], ctx.gpr[31]};
aot_mem.aot_store32_block(ctx.gpr[29] + static_cast<std::uint32_t>(84), tier2_words); }
tier2_gpr_29 = (tier2_gpr_29 + static_cast<std::uint32_t>(-128));
{ const std::uint32_t tier2_words[9]{ctx.gpr[16], tier2_gpr_17, ctx.gpr[18], tier2_gpr_19, ctx.gpr[20], ctx.gpr[21], ctx.gpr[22], ctx.gpr[23], ctx.gpr[31]};
aot_mem.aot_store32_block(tier2_gpr_29 + static_cast<std::uint32_t>(84), tier2_words); }
ctx.gpr[16] = (ctx.gpr[8] | 0u);
ctx.gpr[17] = (ctx.gpr[7] | 0u);
ctx.gpr[18] = (ctx.gpr[6] | 0u);
ctx.gpr[21] = (ctx.gpr[4] | 0u);
ctx.gpr[22] = (ctx.gpr[29] + static_cast<std::uint32_t>(16));
ctx.gpr[23] = (ctx.gpr[29] + static_cast<std::uint32_t>(32));
ctx.gpr[4] = (ctx.gpr[18] | 0u);
ctx.gpr[6] = (ctx.gpr[29] | 0u);
tier2_gpr_17 = (ctx.gpr[7] | 0u);
ctx.gpr[18] = (tier2_gpr_6 | 0u);
ctx.gpr[21] = (tier2_gpr_4 | 0u);
ctx.gpr[22] = (tier2_gpr_29 + static_cast<std::uint32_t>(16));
ctx.gpr[23] = (tier2_gpr_29 + static_cast<std::uint32_t>(32));
tier2_gpr_4 = (ctx.gpr[18] | 0u);
tier2_gpr_6 = (tier2_gpr_29 | 0u);
ctx.gpr[7] = (ctx.gpr[22] | 0u);
ctx.gpr[31] = (0x088B1940u);
ctx.gpr[8] = (ctx.gpr[23] | 0u);
if (rt.invoke_chained_direct<&recomp_unit_0163_entry, 163u, 44u, 0x08A9061Cu>(ctx, &aot_mem) && ctx.pc == 0x088B1940u) goto SB_L_088B1940;
if (([&]() { TIER2_GPR_SYNC_OUT(); const bool tier2_same_ = (rt.invoke_chained_direct<&recomp_unit_0163_entry, 163u, 44u, 0x08A9061Cu>(ctx, &aot_mem)); if (tier2_same_) TIER2_GPR_SYNC_IN(); else tier2_gpr_shadow_valid = false; return tier2_same_; }()) && ctx.pc == 0x088B1940u) goto SB_L_088B1940;
TIER2_SB_RETURN();
SB_L_088B1940:
ctx.gpr[20] = (ctx.gpr[29] + static_cast<std::uint32_t>(48));
ctx.gpr[19] = (ctx.gpr[29] + static_cast<std::uint32_t>(64));
ctx.gpr[10] = (ctx.gpr[29] + static_cast<std::uint32_t>(80));
ctx.gpr[4] = (ctx.gpr[21] | 0u);
ctx.gpr[5] = (ctx.gpr[29] | 0u);
ctx.gpr[6] = (ctx.gpr[22] | 0u);
ctx.gpr[20] = (tier2_gpr_29 + static_cast<std::uint32_t>(48));
tier2_gpr_19 = (tier2_gpr_29 + static_cast<std::uint32_t>(64));
ctx.gpr[10] = (tier2_gpr_29 + static_cast<std::uint32_t>(80));
tier2_gpr_4 = (ctx.gpr[21] | 0u);
tier2_gpr_5 = (tier2_gpr_29 | 0u);
tier2_gpr_6 = (ctx.gpr[22] | 0u);
ctx.gpr[7] = (ctx.gpr[23] | 0u);
ctx.gpr[8] = (ctx.gpr[20] | 0u);
ctx.gpr[31] = (0x088B1968u);
ctx.gpr[9] = (ctx.gpr[19] | 0u);
ctx.gpr[9] = (tier2_gpr_19 | 0u);
goto SB_L_088B1B74;
SB_L_088B1B74:
{ const std::uint32_t vfpu_address = ctx.gpr[4] + static_cast<std::uint32_t>(0);
{ const std::uint32_t vfpu_address = tier2_gpr_4 + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -344,7 +363,7 @@ SB_L_088B1B74:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<1u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[4] + static_cast<std::uint32_t>(16);
{ const std::uint32_t vfpu_address = tier2_gpr_4 + static_cast<std::uint32_t>(16);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -352,7 +371,7 @@ SB_L_088B1B74:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<2u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(0);
{ const std::uint32_t vfpu_address = tier2_gpr_5 + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -360,7 +379,7 @@ SB_L_088B1B74:
std::bit_cast<float>(tier2_vfpu_words[2]),
std::bit_cast<float>(tier2_vfpu_words[3])};
ctx.write_vfpu_vector_ct<4u, 4u>(vfpu_value); }
{ const std::uint32_t vfpu_address = ctx.gpr[6] + static_cast<std::uint32_t>(0);
{ const std::uint32_t vfpu_address = tier2_gpr_6 + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
std::bit_cast<float>(tier2_vfpu_words[0]),
@@ -520,9 +539,9 @@ SB_L_088B1C40:
goto TIER2_LOCAL_DISPATCH_U0043;
SB_L_088B3FCC:
ctx.gpr[18] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(68)));
ctx.gpr[18] = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(68)));
ctx.gpr[18] = (ctx.gpr[18] + ctx.gpr[30]);
ctx.gpr[5] = (ctx.gpr[29] + static_cast<std::uint32_t>(16));
tier2_gpr_5 = (tier2_gpr_29 + static_cast<std::uint32_t>(16));
{ const std::uint32_t vfpu_address = ctx.gpr[18] + static_cast<std::uint32_t>(0);
std::uint32_t tier2_vfpu_words[4]{}; aot_mem.aot_load32_block(vfpu_address, tier2_vfpu_words);
float vfpu_value[4]{
@@ -552,23 +571,23 @@ SB_L_088B3FCC:
}
ctx.write_vfpu_vector_with_destination_prefix_ct<1u, 3u>(vfpu_d); }
{ float vfpu_value[4]{}; ctx.read_vfpu_vector_ct<0u, 4u>(vfpu_value);
const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(0);
const std::uint32_t vfpu_address = tier2_gpr_5 + static_cast<std::uint32_t>(0);
const std::uint32_t tier2_vfpu_words[4]{std::bit_cast<std::uint32_t>(vfpu_value[0]), std::bit_cast<std::uint32_t>(vfpu_value[1]), std::bit_cast<std::uint32_t>(vfpu_value[2]), std::bit_cast<std::uint32_t>(vfpu_value[3])};
aot_mem.aot_store32_block(vfpu_address, tier2_vfpu_words); }
{ float vfpu_value[4]{}; ctx.read_vfpu_vector_ct<1u, 4u>(vfpu_value);
const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(16);
const std::uint32_t vfpu_address = tier2_gpr_5 + static_cast<std::uint32_t>(16);
const std::uint32_t tier2_vfpu_words[4]{std::bit_cast<std::uint32_t>(vfpu_value[0]), std::bit_cast<std::uint32_t>(vfpu_value[1]), std::bit_cast<std::uint32_t>(vfpu_value[2]), std::bit_cast<std::uint32_t>(vfpu_value[3])};
aot_mem.aot_store32_block(vfpu_address, tier2_vfpu_words); }
ctx.gpr[31] = (0x088B3FFCu);
ctx.gpr[4] = (ctx.gpr[20] | 0u);
tier2_gpr_4 = (ctx.gpr[20] | 0u);
goto SB_L_088B1780;
SB_L_088B4004:
ctx.gpr[17] = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(6))))));
ctx.gpr[4] = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(14))))));
ctx.gpr[4] = (static_cast<std::int32_t>(ctx.gpr[4]) < static_cast<std::int32_t>(ctx.gpr[17]) ? 1u : 0u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
ctx.gpr[16] = (ctx.gpr[17] << 3u);
tier2_gpr_17 = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(6))))));
tier2_gpr_4 = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(14))))));
tier2_gpr_4 = (static_cast<std::int32_t>(tier2_gpr_4) < static_cast<std::int32_t>(tier2_gpr_17) ? 1u : 0u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
ctx.gpr[16] = (tier2_gpr_17 << 3u);
if (branch_taken) {
goto SB_L_088B40F4;
}
@@ -585,26 +604,26 @@ SB_L_088B4018:
}
SB_L_088B4020:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[4] = (ctx.gpr[4] + ctx.gpr[16]);
ctx.gpr[4] = (aot_mem.aot_load8(ctx.gpr[4] + static_cast<std::uint32_t>(6)));
ctx.gpr[4] = (ctx.gpr[4] & 255u);
ctx.gpr[5] = (ctx.gpr[4] ^ 7u);
ctx.gpr[5] = (ctx.gpr[5] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[6] = (ctx.gpr[4] ^ 8u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 16u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 31u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[4] = (ctx.gpr[4] ^ 12u);
ctx.gpr[4] = (ctx.gpr[4] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[4] = (ctx.gpr[5] | ctx.gpr[4]);
ctx.gpr[4] = (ctx.gpr[4] & 255u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_4 = (tier2_gpr_4 + ctx.gpr[16]);
tier2_gpr_4 = (aot_mem.aot_load8(tier2_gpr_4 + static_cast<std::uint32_t>(6)));
tier2_gpr_4 = (tier2_gpr_4 & 255u);
tier2_gpr_5 = (tier2_gpr_4 ^ 7u);
tier2_gpr_5 = (tier2_gpr_5 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_6 = (tier2_gpr_4 ^ 8u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 16u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 31u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_4 = (tier2_gpr_4 ^ 12u);
tier2_gpr_4 = (tier2_gpr_4 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_4 = (tier2_gpr_5 | tier2_gpr_4);
tier2_gpr_4 = (tier2_gpr_4 & 255u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
// nop
if (branch_taken) {
goto SB_L_088B40E0;
@@ -622,23 +641,23 @@ SB_L_088B4074:
}
SB_L_088B407C:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[4] = (ctx.gpr[4] + ctx.gpr[16]);
ctx.gpr[4] = (aot_mem.aot_load8(ctx.gpr[4] + static_cast<std::uint32_t>(6)));
ctx.gpr[4] = (ctx.gpr[4] & 255u);
ctx.gpr[5] = (ctx.gpr[4] ^ 8u);
ctx.gpr[5] = (ctx.gpr[5] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[6] = (ctx.gpr[4] ^ 16u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 31u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[4] = (ctx.gpr[4] ^ 12u);
ctx.gpr[4] = (ctx.gpr[4] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[4] = (ctx.gpr[5] | ctx.gpr[4]);
ctx.gpr[4] = (ctx.gpr[4] & 255u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_4 = (tier2_gpr_4 + ctx.gpr[16]);
tier2_gpr_4 = (aot_mem.aot_load8(tier2_gpr_4 + static_cast<std::uint32_t>(6)));
tier2_gpr_4 = (tier2_gpr_4 & 255u);
tier2_gpr_5 = (tier2_gpr_4 ^ 8u);
tier2_gpr_5 = (tier2_gpr_5 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_6 = (tier2_gpr_4 ^ 16u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 31u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_4 = (tier2_gpr_4 ^ 12u);
tier2_gpr_4 = (tier2_gpr_4 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_4 = (tier2_gpr_5 | tier2_gpr_4);
tier2_gpr_4 = (tier2_gpr_4 & 255u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
// nop
if (branch_taken) {
goto SB_L_088B40E0;
@@ -647,13 +666,13 @@ SB_L_088B407C:
}
SB_L_088B40C4:
ctx.gpr[5] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(72)));
ctx.gpr[6] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[6] = (ctx.gpr[6] + ctx.gpr[16]);
ctx.gpr[4] = (ctx.gpr[20] | 0u);
tier2_gpr_5 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(72)));
tier2_gpr_6 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_6 = (tier2_gpr_6 + ctx.gpr[16]);
tier2_gpr_4 = (ctx.gpr[20] | 0u);
ctx.gpr[7] = (ctx.gpr[23] | 0u);
ctx.gpr[31] = (0x088B40E0u);
ctx.gpr[8] = (ctx.gpr[29] | 0u);
ctx.gpr[8] = (tier2_gpr_29 | 0u);
if (tier2_return_depth < kTier2ReturnCapacity) {
if (!rt.tier2_enter_fused_transfer<43u, 0x088B18ECu>(ctx)) {
ctx.pc = 0x088B18ECu;
@@ -667,14 +686,14 @@ SB_L_088B40C4:
goto SB_L_088B18EC;
}
++tier2_stats.fallbacks;
if (rt.invoke_chained_direct<&recomp_unit_0043_entry, 43u, 210u, 0x088B18ECu>(ctx, &aot_mem) && ctx.pc == 0x088B40E0u) goto SB_L_088B40E0;
if (([&]() { TIER2_GPR_SYNC_OUT(); const bool tier2_same_ = (rt.invoke_chained_direct<&recomp_unit_0043_entry, 43u, 210u, 0x088B18ECu>(ctx, &aot_mem)); if (tier2_same_) TIER2_GPR_SYNC_IN(); else tier2_gpr_shadow_valid = false; return tier2_same_; }()) && ctx.pc == 0x088B40E0u) goto SB_L_088B40E0;
TIER2_SB_RETURN();
SB_L_088B40E0:
ctx.gpr[17] = (ctx.gpr[17] + static_cast<std::uint32_t>(1));
ctx.gpr[4] = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(14))))));
ctx.gpr[4] = (static_cast<std::int32_t>(ctx.gpr[4]) < static_cast<std::int32_t>(ctx.gpr[17]) ? 1u : 0u);
{ const bool branch_taken = ctx.gpr[4] == 0u;
tier2_gpr_17 = (tier2_gpr_17 + static_cast<std::uint32_t>(1));
tier2_gpr_4 = (static_cast<std::uint32_t>(static_cast<std::int32_t>(static_cast<std::int16_t>(aot_mem.aot_load16(ctx.gpr[18] + static_cast<std::uint32_t>(14))))));
tier2_gpr_4 = (static_cast<std::int32_t>(tier2_gpr_4) < static_cast<std::int32_t>(tier2_gpr_17) ? 1u : 0u);
{ const bool branch_taken = tier2_gpr_4 == 0u;
ctx.gpr[16] = (ctx.gpr[16] + static_cast<std::uint32_t>(8));
if (branch_taken) {
goto SB_L_088B4018;
@@ -683,16 +702,16 @@ SB_L_088B40E0:
}
SB_L_088B40F4:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(48)));
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(48)));
goto SB_L_088B40F8;
SB_L_088B40F8:
ctx.gpr[4] = (ctx.gpr[4] + static_cast<std::uint32_t>(1));
ctx.gpr[5] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(52)));
tier2_gpr_4 = (tier2_gpr_4 + static_cast<std::uint32_t>(1));
tier2_gpr_5 = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(52)));
ctx.gpr[30] = (ctx.gpr[30] + static_cast<std::uint32_t>(16));
ctx.gpr[5] = (static_cast<std::int32_t>(ctx.gpr[4]) < static_cast<std::int32_t>(ctx.gpr[5]) ? 1u : 0u);
{ const bool branch_taken = ctx.gpr[5] != 0u;
aot_mem.aot_store32(ctx.gpr[29] + static_cast<std::uint32_t>(48), ctx.gpr[4]);
tier2_gpr_5 = (static_cast<std::int32_t>(tier2_gpr_4) < static_cast<std::int32_t>(tier2_gpr_5) ? 1u : 0u);
{ const bool branch_taken = tier2_gpr_5 != 0u;
aot_mem.aot_store32(tier2_gpr_29 + static_cast<std::uint32_t>(48), tier2_gpr_4);
if (branch_taken) {
if (!rt.tier2_enter_fused_transfer<43u, 0x088B3FCCu>(ctx)) {
ctx.pc = 0x088B3FCCu;
@@ -716,10 +735,10 @@ SB_L_088B4110:
SB_L_088B4118:
ctx.gpr[16] = (0u | 0u);
ctx.gpr[4] = (aot_mem.aot_load16(ctx.gpr[19] + static_cast<std::uint32_t>(50)));
ctx.gpr[4] = (static_cast<std::int32_t>(ctx.gpr[16]) < static_cast<std::int32_t>(ctx.gpr[4]) ? 1u : 0u);
{ const bool branch_taken = ctx.gpr[4] == 0u;
ctx.gpr[17] = (0u | 0u);
tier2_gpr_4 = (aot_mem.aot_load16(tier2_gpr_19 + static_cast<std::uint32_t>(50)));
tier2_gpr_4 = (static_cast<std::int32_t>(ctx.gpr[16]) < static_cast<std::int32_t>(tier2_gpr_4) ? 1u : 0u);
{ const bool branch_taken = tier2_gpr_4 == 0u;
tier2_gpr_17 = (0u | 0u);
if (branch_taken) {
goto SB_L_088B4208;
}
@@ -736,26 +755,26 @@ SB_L_088B412C:
}
SB_L_088B4134:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[4] = (ctx.gpr[4] + ctx.gpr[17]);
ctx.gpr[4] = (aot_mem.aot_load8(ctx.gpr[4] + static_cast<std::uint32_t>(6)));
ctx.gpr[4] = (ctx.gpr[4] & 255u);
ctx.gpr[5] = (ctx.gpr[4] ^ 7u);
ctx.gpr[5] = (ctx.gpr[5] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[6] = (ctx.gpr[4] ^ 8u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 16u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 31u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[4] = (ctx.gpr[4] ^ 12u);
ctx.gpr[4] = (ctx.gpr[4] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[4] = (ctx.gpr[5] | ctx.gpr[4]);
ctx.gpr[4] = (ctx.gpr[4] & 255u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_4 = (tier2_gpr_4 + tier2_gpr_17);
tier2_gpr_4 = (aot_mem.aot_load8(tier2_gpr_4 + static_cast<std::uint32_t>(6)));
tier2_gpr_4 = (tier2_gpr_4 & 255u);
tier2_gpr_5 = (tier2_gpr_4 ^ 7u);
tier2_gpr_5 = (tier2_gpr_5 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_6 = (tier2_gpr_4 ^ 8u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 16u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 31u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_4 = (tier2_gpr_4 ^ 12u);
tier2_gpr_4 = (tier2_gpr_4 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_4 = (tier2_gpr_5 | tier2_gpr_4);
tier2_gpr_4 = (tier2_gpr_4 & 255u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
// nop
if (branch_taken) {
goto SB_L_088B41F4;
@@ -773,23 +792,23 @@ SB_L_088B4188:
}
SB_L_088B4190:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[4] = (ctx.gpr[4] + ctx.gpr[17]);
ctx.gpr[4] = (aot_mem.aot_load8(ctx.gpr[4] + static_cast<std::uint32_t>(6)));
ctx.gpr[4] = (ctx.gpr[4] & 255u);
ctx.gpr[5] = (ctx.gpr[4] ^ 8u);
ctx.gpr[5] = (ctx.gpr[5] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[6] = (ctx.gpr[4] ^ 16u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[6] = (ctx.gpr[4] ^ 31u);
ctx.gpr[6] = (ctx.gpr[6] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[5] = (ctx.gpr[5] | ctx.gpr[6]);
ctx.gpr[4] = (ctx.gpr[4] ^ 12u);
ctx.gpr[4] = (ctx.gpr[4] < static_cast<std::uint32_t>(1) ? 1u : 0u);
ctx.gpr[4] = (ctx.gpr[5] | ctx.gpr[4]);
ctx.gpr[4] = (ctx.gpr[4] & 255u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_4 = (tier2_gpr_4 + tier2_gpr_17);
tier2_gpr_4 = (aot_mem.aot_load8(tier2_gpr_4 + static_cast<std::uint32_t>(6)));
tier2_gpr_4 = (tier2_gpr_4 & 255u);
tier2_gpr_5 = (tier2_gpr_4 ^ 8u);
tier2_gpr_5 = (tier2_gpr_5 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_6 = (tier2_gpr_4 ^ 16u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_6 = (tier2_gpr_4 ^ 31u);
tier2_gpr_6 = (tier2_gpr_6 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_5 = (tier2_gpr_5 | tier2_gpr_6);
tier2_gpr_4 = (tier2_gpr_4 ^ 12u);
tier2_gpr_4 = (tier2_gpr_4 < static_cast<std::uint32_t>(1) ? 1u : 0u);
tier2_gpr_4 = (tier2_gpr_5 | tier2_gpr_4);
tier2_gpr_4 = (tier2_gpr_4 & 255u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
// nop
if (branch_taken) {
goto SB_L_088B41F4;
@@ -798,13 +817,13 @@ SB_L_088B4190:
}
SB_L_088B41D8:
ctx.gpr[5] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(72)));
ctx.gpr[6] = (aot_mem.aot_load32(ctx.gpr[19] + static_cast<std::uint32_t>(76)));
ctx.gpr[6] = (ctx.gpr[6] + ctx.gpr[17]);
ctx.gpr[4] = (ctx.gpr[20] | 0u);
tier2_gpr_5 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(72)));
tier2_gpr_6 = (aot_mem.aot_load32(tier2_gpr_19 + static_cast<std::uint32_t>(76)));
tier2_gpr_6 = (tier2_gpr_6 + tier2_gpr_17);
tier2_gpr_4 = (ctx.gpr[20] | 0u);
ctx.gpr[7] = (ctx.gpr[23] | 0u);
ctx.gpr[31] = (0x088B41F4u);
ctx.gpr[8] = (ctx.gpr[29] | 0u);
ctx.gpr[8] = (tier2_gpr_29 | 0u);
if (tier2_return_depth < kTier2ReturnCapacity) {
if (!rt.tier2_enter_fused_transfer<43u, 0x088B18ECu>(ctx)) {
ctx.pc = 0x088B18ECu;
@@ -818,15 +837,15 @@ SB_L_088B41D8:
goto SB_L_088B18EC;
}
++tier2_stats.fallbacks;
if (rt.invoke_chained_direct<&recomp_unit_0043_entry, 43u, 210u, 0x088B18ECu>(ctx, &aot_mem) && ctx.pc == 0x088B41F4u) goto SB_L_088B41F4;
if (([&]() { TIER2_GPR_SYNC_OUT(); const bool tier2_same_ = (rt.invoke_chained_direct<&recomp_unit_0043_entry, 43u, 210u, 0x088B18ECu>(ctx, &aot_mem)); if (tier2_same_) TIER2_GPR_SYNC_IN(); else tier2_gpr_shadow_valid = false; return tier2_same_; }()) && ctx.pc == 0x088B41F4u) goto SB_L_088B41F4;
TIER2_SB_RETURN();
SB_L_088B41F4:
ctx.gpr[16] = (ctx.gpr[16] + static_cast<std::uint32_t>(1));
ctx.gpr[4] = (aot_mem.aot_load16(ctx.gpr[19] + static_cast<std::uint32_t>(50)));
ctx.gpr[4] = (static_cast<std::int32_t>(ctx.gpr[16]) < static_cast<std::int32_t>(ctx.gpr[4]) ? 1u : 0u);
{ const bool branch_taken = ctx.gpr[4] != 0u;
ctx.gpr[17] = (ctx.gpr[17] + static_cast<std::uint32_t>(8));
tier2_gpr_4 = (aot_mem.aot_load16(tier2_gpr_19 + static_cast<std::uint32_t>(50)));
tier2_gpr_4 = (static_cast<std::int32_t>(ctx.gpr[16]) < static_cast<std::int32_t>(tier2_gpr_4) ? 1u : 0u);
{ const bool branch_taken = tier2_gpr_4 != 0u;
tier2_gpr_17 = (tier2_gpr_17 + static_cast<std::uint32_t>(8));
if (branch_taken) {
goto SB_L_088B412C;
}
@@ -834,12 +853,12 @@ SB_L_088B41F4:
}
SB_L_088B4208:
ctx.gpr[4] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(56)));
tier2_gpr_4 = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(56)));
goto SB_L_088B420C;
SB_L_088B420C:
ctx.fpr[12] = std::bit_cast<float>(aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(0)));
ctx.fpr[13] = std::bit_cast<float>(aot_mem.aot_load32(ctx.gpr[4] + static_cast<std::uint32_t>(0)));
ctx.fpr[12] = std::bit_cast<float>(aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(0)));
ctx.fpr[13] = std::bit_cast<float>(aot_mem.aot_load32(tier2_gpr_4 + static_cast<std::uint32_t>(0)));
ctx.fcr31 = (ctx.fcr31 & ~0x00800000u) | (((ctx.fpr[12] < ctx.fpr[13])) ? 0x00800000u : 0u);
// nop
{ const bool branch_taken = !((ctx.fcr31 & 0x00800000u) != 0u);
@@ -851,8 +870,8 @@ SB_L_088B420C:
}
SB_L_088B4224:
ctx.fpr[12] = std::bit_cast<float>(aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(0)));
aot_mem.aot_store32(ctx.gpr[4] + static_cast<std::uint32_t>(0), std::bit_cast<std::uint32_t>(ctx.fpr[12]));
ctx.fpr[12] = std::bit_cast<float>(aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(0)));
aot_mem.aot_store32(tier2_gpr_4 + static_cast<std::uint32_t>(0), std::bit_cast<std::uint32_t>(ctx.fpr[12]));
{ const bool branch_taken = 0u == 0u;
ctx.gpr[2] = (0u | 1u);
if (branch_taken) {
@@ -867,12 +886,12 @@ SB_L_088B4234:
SB_L_088B4238:
{ std::uint32_t tier2_words[11]{};
if (aot_mem.aot_try_load32_block(ctx.gpr[29] + static_cast<std::uint32_t>(60), tier2_words)) {
if (aot_mem.aot_try_load32_block(tier2_gpr_29 + static_cast<std::uint32_t>(60), tier2_words)) {
ctx.fpr[20] = std::bit_cast<float>(tier2_words[0]);
ctx.gpr[16] = tier2_words[1];
ctx.gpr[17] = tier2_words[2];
tier2_gpr_17 = tier2_words[2];
ctx.gpr[18] = tier2_words[3];
ctx.gpr[19] = tier2_words[4];
tier2_gpr_19 = tier2_words[4];
ctx.gpr[20] = tier2_words[5];
ctx.gpr[21] = tier2_words[6];
ctx.gpr[22] = tier2_words[7];
@@ -880,25 +899,28 @@ SB_L_088B4238:
ctx.gpr[30] = tier2_words[9];
ctx.gpr[31] = tier2_words[10];
} else {
ctx.fpr[20] = std::bit_cast<float>(aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(60)));
ctx.gpr[16] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(64)));
ctx.gpr[17] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(68)));
ctx.gpr[18] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(72)));
ctx.gpr[19] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(76)));
ctx.gpr[20] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(80)));
ctx.gpr[21] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(84)));
ctx.gpr[22] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(88)));
ctx.gpr[23] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(92)));
ctx.gpr[30] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(96)));
ctx.gpr[31] = (aot_mem.aot_load32(ctx.gpr[29] + static_cast<std::uint32_t>(100)));
ctx.fpr[20] = std::bit_cast<float>(aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(60)));
ctx.gpr[16] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(64)));
tier2_gpr_17 = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(68)));
ctx.gpr[18] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(72)));
tier2_gpr_19 = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(76)));
ctx.gpr[20] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(80)));
ctx.gpr[21] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(84)));
ctx.gpr[22] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(88)));
ctx.gpr[23] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(92)));
ctx.gpr[30] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(96)));
ctx.gpr[31] = (aot_mem.aot_load32(tier2_gpr_29 + static_cast<std::uint32_t>(100)));
} }
jump_target = ctx.gpr[31];
ctx.gpr[29] = (ctx.gpr[29] + static_cast<std::uint32_t>(112));
tier2_gpr_29 = (tier2_gpr_29 + static_cast<std::uint32_t>(112));
local_pc = jump_target;
goto TIER2_LOCAL_DISPATCH_U0044;
// TIER2_GPR_BODY_END
#undef TIER2_SB_RETURN
#undef TIER2_GPR_BEFORE_COLD
#undef TIER2_GPR_SYNC_IN
#undef TIER2_GPR_SYNC_OUT
}
} // namespace vcs
profiles/vcs/host/vcs_tier2_cluster_entity.cpp
+353 -319
View File
File diff suppressed because it is too large Load Diff
profiles/vcs/host/vcs_tier2_cluster_geometry.cpp
+89 -53
View File
@@ -1,8 +1,9 @@
// AUTO-GENERATED by profiles/vcs/tools/build_tier2_superblocks.py.
// Tier-2 SUPERBLOCK V3 DATAFLOW cluster: geometry
// Tier-2 SUPERBLOCK V4 150FPS cluster: geometry
#include "vcs_tier2_superblocks.hpp"
#include "psprecomp/runtime.hpp"
#include "generated_units.hpp"
#include "vcs_fast_paths.hpp"
#include <cstdint>
@@ -27,8 +28,19 @@ void tier2_superblock_geometry(psprecomp::Runtime &rt,
std::uint32_t local_pc = 0u;
std::uint32_t local_transfers = 0u;
std::uint32_t entry_id = 0u;
bool tier2_gpr_shadow_valid = true;
#define TIER2_GPR_SYNC_OUT() do {} while (false)
#define TIER2_GPR_SYNC_IN() do {} while (false)
#define TIER2_GPR_BEFORE_COLD() do { tier2_gpr_shadow_valid = false; } while (false)
#define TIER2_SB_RETURN() do { tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)rt.tier2_complete_fused_transfers(ctx, tier2_tail_count_); --tier2_return_depth; (void)rt.tier2_complete_fused_transfers(ctx, 1u); } if (tier2_pending_transfers != 0u) { (void)rt.tier2_complete_fused_transfers(ctx, tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
auto tier2_complete_shadow = [&](std::uint32_t count) -> bool {
TIER2_GPR_SYNC_OUT();
const bool same = rt.tier2_complete_fused_transfers(ctx, count);
if (!same) tier2_gpr_shadow_valid = false;
return same;
};
#define TIER2_SB_RETURN() do { TIER2_GPR_SYNC_OUT(); tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)tier2_complete_shadow(tier2_tail_count_); --tier2_return_depth; (void)tier2_complete_shadow(1u); } if (tier2_pending_transfers != 0u) { (void)tier2_complete_shadow(tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
goto TIER2_ENTRY_DISPATCH;
@@ -443,11 +455,11 @@ TIER2_LOCAL_DISPATCH_U0084:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -539,11 +551,11 @@ TIER2_LOCAL_DISPATCH_U0085:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -886,6 +898,7 @@ TIER2_ENTRY_DISPATCH:
TIER2_SB_RETURN();
}
// TIER2_GPR_BODY_BEGIN
SB_L_08955444:
ctx.gpr[9] = (aot_mem.aot_load32(ctx.gpr[4] + static_cast<std::uint32_t>(11768)));
ctx.set_vfpu_scalar_bits_ct<12u>(aot_mem.aot_load32(ctx.gpr[5] + static_cast<std::uint32_t>(4)));
@@ -905,10 +918,14 @@ SB_L_08955444:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -1043,10 +1060,14 @@ SB_L_0895554C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -1638,25 +1659,13 @@ SB_L_08957500:
{ float vfpu_s[16]{}, vfpu_t[16]{}, vfpu_d[16]{};
ctx.read_vfpu_matrix_ct<36u, 4u>(vfpu_s);
ctx.read_vfpu_matrix_ct<24u, 4u>(vfpu_t);
for (std::uint32_t a = 0; a < 4u; ++a) {
for (std::uint32_t b = 0; b < 4u; ++b) {
float sum = 0.0f;
for (std::uint32_t c = 0; c < 4u; ++c) sum += vfpu_s[b * 4u + c] * vfpu_t[a * 4u + c];
vfpu_d[a * 4u + b] = sum;
}
}
psprecomp::vcs_tier2_mat4_mul_ordered(vfpu_s, vfpu_t, vfpu_d);
ctx.write_vfpu_matrix_ct<20u, 4u>(vfpu_d);
ctx.eat_vfpu_prefixes(); }
{ float vfpu_s[16]{}, vfpu_t[16]{}, vfpu_d[16]{};
ctx.read_vfpu_matrix_ct<40u, 4u>(vfpu_s);
ctx.read_vfpu_matrix_ct<24u, 4u>(vfpu_t);
for (std::uint32_t a = 0; a < 4u; ++a) {
for (std::uint32_t b = 0; b < 4u; ++b) {
float sum = 0.0f;
for (std::uint32_t c = 0; c < 4u; ++c) sum += vfpu_s[b * 4u + c] * vfpu_t[a * 4u + c];
vfpu_d[a * 4u + b] = sum;
}
}
psprecomp::vcs_tier2_mat4_mul_ordered(vfpu_s, vfpu_t, vfpu_d);
ctx.write_vfpu_matrix_ct<16u, 4u>(vfpu_d);
ctx.eat_vfpu_prefixes(); }
ctx.execute_vfpu_vcmp_ct<52u, 31u, 3u, 6u>();
@@ -2377,10 +2386,14 @@ SB_L_08959118:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -3434,10 +3447,14 @@ SB_L_089598B0:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -3550,10 +3567,14 @@ SB_L_08959930:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -4071,10 +4092,14 @@ SB_L_08959BF0:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -4265,10 +4290,14 @@ SB_L_08959CD8:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -4793,10 +4822,14 @@ SB_L_0895A03C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -6900,9 +6933,12 @@ SB_L_0895B33C:
ctx.gpr[29] = (ctx.gpr[29] + static_cast<std::uint32_t>(144));
local_pc = jump_target;
goto TIER2_LOCAL_DISPATCH_U0085;
// TIER2_GPR_BODY_END
#undef TIER2_SB_RETURN
#undef TIER2_GPR_BEFORE_COLD
#undef TIER2_GPR_SYNC_IN
#undef TIER2_GPR_SYNC_OUT
}
} // namespace vcs
profiles/vcs/host/vcs_tier2_cluster_matrix.cpp
+63 -39
View File
@@ -1,8 +1,9 @@
// AUTO-GENERATED by profiles/vcs/tools/build_tier2_superblocks.py.
// Tier-2 SUPERBLOCK V3 DATAFLOW cluster: matrix
// Tier-2 SUPERBLOCK V4 150FPS cluster: matrix
#include "vcs_tier2_superblocks.hpp"
#include "psprecomp/runtime.hpp"
#include "generated_units.hpp"
#include "vcs_fast_paths.hpp"
#include <cstdint>
@@ -27,8 +28,19 @@ void tier2_superblock_matrix(psprecomp::Runtime &rt,
std::uint32_t local_pc = 0u;
std::uint32_t local_transfers = 0u;
std::uint32_t entry_id = 0u;
bool tier2_gpr_shadow_valid = true;
#define TIER2_GPR_SYNC_OUT() do {} while (false)
#define TIER2_GPR_SYNC_IN() do {} while (false)
#define TIER2_GPR_BEFORE_COLD() do { tier2_gpr_shadow_valid = false; } while (false)
#define TIER2_SB_RETURN() do { tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)rt.tier2_complete_fused_transfers(ctx, tier2_tail_count_); --tier2_return_depth; (void)rt.tier2_complete_fused_transfers(ctx, 1u); } if (tier2_pending_transfers != 0u) { (void)rt.tier2_complete_fused_transfers(ctx, tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
auto tier2_complete_shadow = [&](std::uint32_t count) -> bool {
TIER2_GPR_SYNC_OUT();
const bool same = rt.tier2_complete_fused_transfers(ctx, count);
if (!same) tier2_gpr_shadow_valid = false;
return same;
};
#define TIER2_SB_RETURN() do { TIER2_GPR_SYNC_OUT(); tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)tier2_complete_shadow(tier2_tail_count_); --tier2_return_depth; (void)tier2_complete_shadow(1u); } if (tier2_pending_transfers != 0u) { (void)tier2_complete_shadow(tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
goto TIER2_ENTRY_DISPATCH;
@@ -137,11 +149,11 @@ TIER2_LOCAL_DISPATCH_U0044:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -235,6 +247,7 @@ TIER2_ENTRY_DISPATCH:
TIER2_SB_RETURN();
}
// TIER2_GPR_BODY_BEGIN
SB_L_088B4738:
ctx.gpr[29] = (ctx.gpr[29] + static_cast<std::uint32_t>(-352));
{ const std::uint32_t tier2_words[11]{std::bit_cast<std::uint32_t>(ctx.fpr[20]), ctx.gpr[16], ctx.gpr[17], ctx.gpr[18], ctx.gpr[19], ctx.gpr[20], ctx.gpr[21], ctx.gpr[22], ctx.gpr[23], ctx.gpr[30], ctx.gpr[31]};
@@ -314,10 +327,14 @@ SB_L_088B4738:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -340,13 +357,7 @@ SB_L_088B4738:
{ float vfpu_s[16]{}, vfpu_t[16]{}, vfpu_d[16]{};
ctx.read_vfpu_matrix(vfpu_s, 0u, 4u);
ctx.read_vfpu_matrix(vfpu_t, 36u, 4u);
for (std::uint32_t a = 0; a < 4u; ++a) {
for (std::uint32_t b = 0; b < 4u; ++b) {
float sum = 0.0f;
for (std::uint32_t c = 0; c < 4u; ++c) sum += vfpu_s[b * 4u + c] * vfpu_t[a * 4u + c];
vfpu_d[a * 4u + b] = sum;
}
}
psprecomp::vcs_tier2_mat4_mul_ordered(vfpu_s, vfpu_t, vfpu_d);
ctx.write_vfpu_matrix(vfpu_d, 44u, 4u);
ctx.eat_vfpu_prefixes(); }
{ const std::uint32_t vfpu_address = ctx.gpr[5] + static_cast<std::uint32_t>(0);
@@ -393,10 +404,14 @@ SB_L_088B4738:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -457,10 +472,14 @@ SB_L_088B47F0:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -487,13 +506,7 @@ SB_L_088B47F0:
{ float vfpu_s[16]{}, vfpu_t[16]{}, vfpu_d[16]{};
ctx.read_vfpu_matrix(vfpu_s, 36u, 4u);
ctx.read_vfpu_matrix(vfpu_t, 0u, 4u);
for (std::uint32_t a = 0; a < 4u; ++a) {
for (std::uint32_t b = 0; b < 4u; ++b) {
float sum = 0.0f;
for (std::uint32_t c = 0; c < 4u; ++c) sum += vfpu_s[b * 4u + c] * vfpu_t[a * 4u + c];
vfpu_d[a * 4u + b] = sum;
}
}
psprecomp::vcs_tier2_mat4_mul_ordered(vfpu_s, vfpu_t, vfpu_d);
ctx.write_vfpu_matrix(vfpu_d, 48u, 4u);
ctx.eat_vfpu_prefixes(); }
goto SB_L_088B4810;
@@ -1054,10 +1067,14 @@ SB_L_088B4C7C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -1392,10 +1409,14 @@ SB_L_088B4EE4:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -1468,9 +1489,12 @@ SB_L_088B4F74:
ctx.gpr[29] = (ctx.gpr[29] + static_cast<std::uint32_t>(352));
local_pc = jump_target;
goto TIER2_LOCAL_DISPATCH_U0044;
// TIER2_GPR_BODY_END
#undef TIER2_SB_RETURN
#undef TIER2_GPR_BEFORE_COLD
#undef TIER2_GPR_SYNC_IN
#undef TIER2_GPR_SYNC_OUT
}
} // namespace vcs
profiles/vcs/host/vcs_tier2_cluster_physics.cpp
+45 -17
View File
@@ -1,8 +1,9 @@
// AUTO-GENERATED by profiles/vcs/tools/build_tier2_superblocks.py.
// Tier-2 SUPERBLOCK V3 DATAFLOW cluster: physics
// Tier-2 SUPERBLOCK V4 150FPS cluster: physics
#include "vcs_tier2_superblocks.hpp"
#include "psprecomp/runtime.hpp"
#include "generated_units.hpp"
#include "vcs_fast_paths.hpp"
#include <cstdint>
@@ -27,8 +28,19 @@ void tier2_superblock_physics(psprecomp::Runtime &rt,
std::uint32_t local_pc = 0u;
std::uint32_t local_transfers = 0u;
std::uint32_t entry_id = 0u;
bool tier2_gpr_shadow_valid = true;
#define TIER2_GPR_SYNC_OUT() do {} while (false)
#define TIER2_GPR_SYNC_IN() do {} while (false)
#define TIER2_GPR_BEFORE_COLD() do { tier2_gpr_shadow_valid = false; } while (false)
#define TIER2_SB_RETURN() do { tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)rt.tier2_complete_fused_transfers(ctx, tier2_tail_count_); --tier2_return_depth; (void)rt.tier2_complete_fused_transfers(ctx, 1u); } if (tier2_pending_transfers != 0u) { (void)rt.tier2_complete_fused_transfers(ctx, tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
auto tier2_complete_shadow = [&](std::uint32_t count) -> bool {
TIER2_GPR_SYNC_OUT();
const bool same = rt.tier2_complete_fused_transfers(ctx, count);
if (!same) tier2_gpr_shadow_valid = false;
return same;
};
#define TIER2_SB_RETURN() do { TIER2_GPR_SYNC_OUT(); tier2_scope.finish(); /* Unwind every logical invoke_chained_direct frame in true LIFO order. Tail frames created inside a fused JAL unwind before that JAL; outer JAL frames are also released after context invalidation. */ while (tier2_return_depth != 0u) { std::uint32_t tier2_base_ = tier2_return_pending_base[tier2_return_depth - 1u]; if (tier2_base_ > tier2_pending_transfers) { ++tier2_stats.fallbacks; tier2_base_ = tier2_pending_transfers; } const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; tier2_pending_transfers = tier2_base_; if (tier2_tail_count_ != 0u) (void)tier2_complete_shadow(tier2_tail_count_); --tier2_return_depth; (void)tier2_complete_shadow(1u); } if (tier2_pending_transfers != 0u) { (void)tier2_complete_shadow(tier2_pending_transfers); tier2_pending_transfers = 0u; } return; } while (false)
goto TIER2_ENTRY_DISPATCH;
@@ -79,11 +91,11 @@ TIER2_LOCAL_DISPATCH_U0129:
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -119,6 +131,7 @@ TIER2_ENTRY_DISPATCH:
TIER2_SB_RETURN();
}
// TIER2_GPR_BODY_BEGIN
SB_L_08A094D8:
{ const bool branch_taken = ctx.gpr[5] == 0u;
ctx.gpr[6] = (ctx.gpr[6] & 1u);
@@ -220,10 +233,14 @@ SB_L_08A09B2C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -262,10 +279,14 @@ SB_L_08A09B2C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -331,10 +352,14 @@ SB_L_08A09C1C:
constexpr std::uint32_t vfpu_input_length = 3u;
for (std::uint32_t i = 0; i < 4u; ++i) vfpu_target[i] = i < vfpu_input_length ? vfpu_target_raw[i] : 0.0f;
if (vfpu_side - 1u >= vfpu_input_length) vfpu_target[vfpu_side - 1u] = 1.0f;
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
if constexpr (vfpu_side == 4u) {
psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);
} else {
for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {
float sum = 0.0f;
for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];
vfpu_result[row] = sum;
}
}
float vfpu_final_row[4]{vfpu_matrix[(vfpu_side - 1u) * 4u + 0u], vfpu_matrix[(vfpu_side - 1u) * 4u + 1u],
vfpu_matrix[(vfpu_side - 1u) * 4u + 2u], vfpu_matrix[(vfpu_side - 1u) * 4u + 3u]};
@@ -414,9 +439,12 @@ SB_L_08A09CBC:
ctx.gpr[31] = (0x08A09CC4u);
ctx.gpr[6] = (ctx.gpr[19] ^ ctx.gpr[5]);
goto SB_L_08A094D8;
// TIER2_GPR_BODY_END
#undef TIER2_SB_RETURN
#undef TIER2_GPR_BEFORE_COLD
#undef TIER2_GPR_SYNC_IN
#undef TIER2_GPR_SYNC_OUT
}
} // namespace vcs
profiles/vcs/host/vcs_tier2_cluster_world.cpp
+1393 -1371
View File
File diff suppressed because it is too large Load Diff
profiles/vcs/scripts/build_release_ninja.bat
+8 -9
View File
@@ -92,7 +92,7 @@ if not exist "%CTEST_EXE%" set "CTEST_EXE=ctest.exe"
set "NINJA_STATUS=[%%f/%%t %%p ^| %%e elapsed ^| %%r running] "
set "BOOTFIX_STAMP=%BUILD%\.vcs_tier2_bootfix_20260816_v1"
set "SUPERBLOCK_STAMP=%BUILD%\.vcs_tier2_superblock_v3_dataflow_20260816"
set "SUPERBLOCK_STAMP=%BUILD%\.vcs_tier2_v4_150fps_buildfix3_20260816"
echo ================================================================
echo VCS - NINJA PERFORMANCE INCREMENTAL BUILD
@@ -104,7 +104,7 @@ echo CMake: %CMAKE_EXE%
echo Ninja: %NINJA_EXE%
echo Ninja workers: %JOBS%
echo cl.exe /MP: OFF ^(Ninja owns compile parallelism^)
echo Generated AOT: O3, cold /Ob0, measured hot /Ob3; Tier2 V3 dataflow clusters O2 /Ob3 /GL-
echo Generated AOT: O3, cold /Ob0, measured hot /Ob3; Tier2 V4 150FPS clusters O2 /Ob3; Geometry /Ob2 + shadow-off compile-guard; /GL- + GPR shadow + SIMD
echo Host/core LTCG: ON
echo AVX2/fast paths: ON
echo ================================================================
@@ -113,7 +113,7 @@ echo [0b/7] Reapplying BOOTFIX-safe Tier-2 transforms (OPT1 semantic transforms
call "%PROFILE%\APPLY_TIER2_EXTREME.bat"
if errorlevel 1 goto :FAIL
echo [0b2/7] Building profile-guided Tier-2 V3 dataflow second layer...
echo [0b2/7] Building profile-guided Tier-2 V4 150FPS second layer...
set "PYTHON3_CMD="
py -3 -c "import sys; raise SystemExit(0 if sys.version_info.major == 3 else 1)" >nul 2>&1
if not errorlevel 1 set "PYTHON3_CMD=py -3"
@@ -128,11 +128,10 @@ if errorlevel 1 goto :FAIL
if exist "%BUILD%" if not exist "%SUPERBLOCK_STAMP%" (
echo.
echo [0c-super/7] Tier2 SUPERBLOCK V3 DATAFLOW - invalidating measured hook/cluster objects once...
for %%U in (0043 0044 0084 0085 0086 0129 0154 0155 0157 0158) do del /s /q "%BUILD%\*generated_unit_%%U*.obj" >nul 2>&1
del /s /q "%BUILD%\*vcs_tier2_superblocks*.obj" >nul 2>&1
del /s /q "%BUILD%\*vcs_tier2_cluster_*.obj" >nul 2>&1
del /s /q "%BUILD%\*vcs_profile*.obj" >nul 2>&1
echo [0c-super/7] Tier2 V4 BUILDFIX3 - invalidating Geometry + runtime log once...
rem BUILDFIX3 only changes Geometry codegen/shadow policy and runtime metadata.
rem Keep every already-valid V4 object so Ninja does not repeat the expensive build.
del /s /q "%BUILD%\*vcs_tier2_cluster_geometry*.obj" >nul 2>&1
del /s /q "%BUILD%\*vcs_runtime_log*.obj" >nul 2>&1
)
@@ -170,7 +169,7 @@ echo [2/7] Building VCSNative with Ninja...
"%CMAKE_EXE%" --build "%BUILD%" --parallel %JOBS% --target VCSNative
if errorlevel 1 goto :FAIL
>"%BOOTFIX_STAMP%" echo VCS Tier2 BOOTFIX 2026-08-16 v1
>"%SUPERBLOCK_STAMP%" echo VCS Tier2 SUPERBLOCK V3 DATAFLOW 2026-08-16
>"%SUPERBLOCK_STAMP%" echo VCS Tier2 V4 BUILDFIX3 2026-08-16
echo.
echo [2b/7] Building tests and DX12 probes...
profiles/vcs/tools/build_tier2_superblocks.py
+128 -16
View File
@@ -1,7 +1,7 @@
#!/usr/bin/env python3
"""Generate VCS Tier-2 SUPERBLOCK V3 dataflow multi-cluster second-layer AOT.
"""Generate VCS Tier-2 SUPERBLOCK V4 150FPS multi-cluster second-layer AOT.
V3 is profile-guided and intentionally keeps the original generated corpus as
V4 is profile-guided and intentionally keeps the original generated corpus as
its semantic fallback. It extracts only measured hot control-flow closures,
fuses selected cross-unit direct calls/tails inside those closures, and leaves
all cold/external paths in the original AOT units.
@@ -137,7 +137,7 @@ CLUSTERS: List[Cluster] = [
# Tier-2 V3 dataflow/memory lowering. V2 proved that eliminating dispatch
# Tier-2 V4 dataflow/memory lowering inherited from V3. V2 proved that eliminating dispatch
# alone is not enough: the hot clusters still spend most of their time in the
# translated body. These transforms deliberately target memory-access runs
# where semantics can be preserved without keeping guest registers dirty
@@ -283,6 +283,102 @@ def _batch_simple_store_runs(text: str) -> tuple[str, int, int]:
return '\n'.join(out) + ('\n' if text.endswith('\n') else ''), runs, words
MAT4_MUL_RE = re.compile(
r'for \(std::uint32_t a = 0; a < 4u; \+\+a\) \{\s*'
r'for \(std::uint32_t b = 0; b < 4u; \+\+b\) \{\s*'
r'float sum = 0\.0f;\s*'
r'for \(std::uint32_t c = 0; c < 4u; \+\+c\) sum \+= vfpu_s\[b \* 4u \+ c\] \* vfpu_t\[a \* 4u \+ c\];\s*'
r'vfpu_d\[a \* 4u \+ b\] = sum;\s*'
r'\}\s*\}')
MAT4_VEC_FIRST3_RE = re.compile(
r'for \(std::uint32_t row = 0; row \+ 1u < vfpu_side; \+\+row\) \{\s*'
r'float sum = 0\.0f;\s*'
r'for \(std::uint32_t column = 0; column < vfpu_side; \+\+column\) sum \+= vfpu_matrix\[row \* 4u \+ column\] \* vfpu_target\[column\];\s*'
r'vfpu_result\[row\] = sum;\s*'
r'\}')
def optimize_tier2_simd(text: str) -> tuple[str, Dict[str, int]]:
stats = {'mat4_mul': 0, 'mat4_vec_first3': 0}
def repl_mul(_: re.Match[str]) -> str:
stats['mat4_mul'] += 1
return 'psprecomp::vcs_tier2_mat4_mul_ordered(vfpu_s, vfpu_t, vfpu_d);'
text = MAT4_MUL_RE.sub(repl_mul, text)
def repl_vec(_: re.Match[str]) -> str:
stats['mat4_vec_first3'] += 1
return ('if constexpr (vfpu_side == 4u) {\n'
' psprecomp::vcs_tier2_mat4_vec_first3_ordered(vfpu_matrix, vfpu_target, vfpu_result);\n'
' } else {\n'
' for (std::uint32_t row = 0; row + 1u < vfpu_side; ++row) {\n'
' float sum = 0.0f;\n'
' for (std::uint32_t column = 0; column < vfpu_side; ++column) sum += vfpu_matrix[row * 4u + column] * vfpu_target[column];\n'
' vfpu_result[row] = sum;\n'
' }\n'
' }')
text = MAT4_VEC_FIRST3_RE.sub(repl_vec, text)
return text, stats
_GPR_BODY_BEGIN = '// TIER2_GPR_BODY_BEGIN\n'
_GPR_BODY_END = '// TIER2_GPR_BODY_END\n'
_GPR_DECL_PLACEHOLDER = ' // TIER2_GPR_SHADOW_DECLS\n'
_GPR_DIRECT_BOOL_RE = re.compile(r'rt\.invoke_chained_direct<[^;\n]+?>\(ctx, &aot_mem\)')
_GPR_DYNAMIC_BOOL_RE = re.compile(r'rt\.invoke_chained_call\(ctx, &aot_mem\)')
def optimize_tier2_gpr_shadow(text: str, cluster_key: str) -> tuple[str, Dict[str, object]]:
"""Promote six hot GPRs with explicit synchronization at visibility boundaries."""
begin = text.find(_GPR_BODY_BEGIN)
end = text.find(_GPR_BODY_END)
if begin < 0 or end < 0 or end <= begin:
raise RuntimeError(f'{cluster_key}: GPR body markers missing')
body_start = begin + len(_GPR_BODY_BEGIN)
body = text[body_start:end]
eligible = cluster_key not in {'physics', 'matrix', 'geometry'}
counts = collections.Counter(int(x) for x in re.findall(r'ctx\.gpr\[(\d+)\]', body))
selected = [reg for reg, _ in counts.most_common() if reg != 0][:6] if eligible else []
if selected:
decls = ''.join(f' std::uint32_t tier2_gpr_{r} = ctx.gpr[{r}];\n' for r in selected)
out_body = ' '.join(f'ctx.gpr[{r}] = tier2_gpr_{r};' for r in selected)
in_body = ' '.join(f'tier2_gpr_{r} = ctx.gpr[{r}];' for r in selected)
macros = (decls +
' bool tier2_gpr_shadow_valid = true;\n' +
f'#define TIER2_GPR_SYNC_OUT() do {{ if (tier2_gpr_shadow_valid) {{ {out_body} }} }} while (false)\n' +
f'#define TIER2_GPR_SYNC_IN() do {{ if (tier2_gpr_shadow_valid) {{ {in_body} }} }} while (false)\n' +
'#define TIER2_GPR_BEFORE_COLD() do { TIER2_GPR_SYNC_OUT(); tier2_gpr_shadow_valid = false; } while (false)\n')
else:
macros = (' bool tier2_gpr_shadow_valid = true;\n'
'#define TIER2_GPR_SYNC_OUT() do {} while (false)\n'
'#define TIER2_GPR_SYNC_IN() do {} while (false)\n'
'#define TIER2_GPR_BEFORE_COLD() do { tier2_gpr_shadow_valid = false; } while (false)\n')
if _GPR_DECL_PLACEHOLDER not in text:
raise RuntimeError(f'{cluster_key}: GPR declaration placeholder missing')
if selected:
def wrap_bool(m: re.Match[str]) -> str:
expr = m.group(0)
return ('([&]() { TIER2_GPR_SYNC_OUT(); const bool tier2_same_ = (' + expr + '); '
'if (tier2_same_) TIER2_GPR_SYNC_IN(); else tier2_gpr_shadow_valid = false; '
'return tier2_same_; }())')
body = _GPR_DIRECT_BOOL_RE.sub(wrap_bool, body)
body = _GPR_DYNAMIC_BOOL_RE.sub(wrap_bool, body)
for r in selected:
body = body.replace(f'ctx.gpr[{r}]', f'tier2_gpr_{r}')
text = text[:body_start] + body + text[end:]
# Insert declarations only after body splicing; doing this earlier shifts
# the saved marker offsets and corrupts the generated CFG.
text = text.replace(_GPR_DECL_PLACEHOLDER, macros, 1)
return text, {
'enabled': int(bool(selected)),
'registers': selected,
'occurrences': sum(counts[r] for r in selected),
}
def optimize_tier2_dataflow(text: str) -> tuple[str, Dict[str, int]]:
stats: Dict[str, int] = {
'vfpu_quad_loads': 0, 'vfpu_quad_stores': 0, 'append32': 0, 'advance32': 0,
@@ -526,7 +622,7 @@ def transform_block(block: str, source_unit: int, selected: Mapping[int, Set[int
stats['cold_exits'] += 1
if entry is not None:
return (f'++tier2_stats.cold_exits; tier2_scope.finish(); '
f'psprecomp::recomp_unit_{source_unit:04d}_entry(rt, ctx, {entry}u, aot_mem); '
f'TIER2_GPR_BEFORE_COLD(); psprecomp::recomp_unit_{source_unit:04d}_entry(rt, ctx, {entry}u, aot_mem); '
f'TIER2_SB_RETURN();')
return (f'++tier2_stats.cold_exits; ctx.pc = 0x{target:08X}u; TIER2_SB_RETURN();')
text = GOTO_RE.sub(goto_repl, text)
@@ -589,11 +685,11 @@ def emit_cluster(cluster: Cluster, selected: Mapping[int, Set[int]],
tier2_pending_transfers - tier2_pending_base;
if (tier2_nested_tail != 0u) {{
tier2_pending_transfers = tier2_pending_base;
if (!rt.tier2_complete_fused_transfers(ctx, tier2_nested_tail))
if (!tier2_complete_shadow(tier2_nested_tail))
tier2_same_context = false;
}}
--tier2_return_depth;
if (!rt.tier2_complete_fused_transfers(ctx, 1u))
if (!tier2_complete_shadow(1u))
tier2_same_context = false;
if (!tier2_same_context) TIER2_SB_RETURN();
goto TIER2_FUSED_RETURN_DISPATCH;
@@ -619,7 +715,7 @@ def emit_cluster(cluster: Cluster, selected: Mapping[int, Set[int]],
for unit, entries in sorted(cold_resume_by_unit.items()):
cases = '\n'.join(
f' case 0x{pc:08X}u: ++tier2_stats.cold_exits; tier2_scope.finish(); '
f'psprecomp::recomp_unit_{unit:04d}_entry(rt, ctx, {entry}u, aot_mem); TIER2_SB_RETURN();'
f'TIER2_GPR_BEFORE_COLD(); psprecomp::recomp_unit_{unit:04d}_entry(rt, ctx, {entry}u, aot_mem); TIER2_SB_RETURN();'
for pc, entry in entries)
cold_resume_sections.append(f''' case {unit}u:
switch (tier2_resume_pc) {{
@@ -629,10 +725,11 @@ def emit_cluster(cluster: Cluster, selected: Mapping[int, Set[int]],
break;''')
cpp = f'''// AUTO-GENERATED by profiles/vcs/tools/build_tier2_superblocks.py.
// Tier-2 SUPERBLOCK V3 DATAFLOW cluster: {cluster.key}
// Tier-2 SUPERBLOCK V4 150FPS cluster: {cluster.key}
#include "vcs_tier2_superblocks.hpp"
#include "psprecomp/runtime.hpp"
#include "generated_units.hpp"
#include "vcs_fast_paths.hpp"
#include <cstdint>
@@ -657,8 +754,17 @@ void {cluster.function}(psprecomp::Runtime &rt,
std::uint32_t local_pc = 0u;
std::uint32_t local_transfers = 0u;
std::uint32_t entry_id = 0u;
// TIER2_GPR_SHADOW_DECLS
auto tier2_complete_shadow = [&](std::uint32_t count) -> bool {{
TIER2_GPR_SYNC_OUT();
const bool same = rt.tier2_complete_fused_transfers(ctx, count);
if (!same) tier2_gpr_shadow_valid = false;
return same;
}};
#define TIER2_SB_RETURN() do {{ \
TIER2_GPR_SYNC_OUT(); \
tier2_scope.finish(); \
/* Unwind every logical invoke_chained_direct frame in true LIFO \
order. Tail frames created inside a fused JAL unwind before that \
@@ -672,12 +778,12 @@ void {cluster.function}(psprecomp::Runtime &rt,
const std::uint32_t tier2_tail_count_ = tier2_pending_transfers - tier2_base_; \
tier2_pending_transfers = tier2_base_; \
if (tier2_tail_count_ != 0u) \
(void)rt.tier2_complete_fused_transfers(ctx, tier2_tail_count_); \
(void)tier2_complete_shadow(tier2_tail_count_); \
--tier2_return_depth; \
(void)rt.tier2_complete_fused_transfers(ctx, 1u); \
(void)tier2_complete_shadow(1u); \
}} \
if (tier2_pending_transfers != 0u) {{ \
(void)rt.tier2_complete_fused_transfers(ctx, tier2_pending_transfers); \
(void)tier2_complete_shadow(tier2_pending_transfers); \
tier2_pending_transfers = 0u; \
}} \
return; \
@@ -708,15 +814,21 @@ TIER2_ENTRY_DISPATCH:
TIER2_SB_RETURN();
}}
{chr(10).join(chunks)}
#undef TIER2_SB_RETURN
// TIER2_GPR_BODY_BEGIN\n{chr(10).join(chunks)}// TIER2_GPR_BODY_END\n\n#undef TIER2_SB_RETURN
#undef TIER2_GPR_BEFORE_COLD
#undef TIER2_GPR_SYNC_IN
#undef TIER2_GPR_SYNC_OUT
}}
}} // namespace vcs
'''
cpp, dataflow_stats = optimize_tier2_dataflow(cpp)
stats.update({f'dataflow_{k}': v for k, v in dataflow_stats.items()})
cpp, simd_stats = optimize_tier2_simd(cpp)
stats.update({f'simd_{k}': v for k, v in simd_stats.items()})
cpp, gpr_stats = optimize_tier2_gpr_shadow(cpp, cluster.key)
stats['gpr_shadow_occurrences'] = int(gpr_stats['occurrences'])
stats['gpr_shadow_registers'] = len(gpr_stats['registers'])
path = host / f'vcs_tier2_cluster_{cluster.key}.cpp'
old = path.read_text(encoding='utf-8') if path.exists() else None
@@ -779,7 +891,7 @@ def main() -> int:
selected_by_cluster[cluster.key] = selected
path, stats = emit_cluster(cluster, selected, sources, host)
generated_paths.append(path)
print(f'Tier2 V3 cluster {cluster.key}: ' + ' '.join(f'{k}={v}' for k, v in sorted(stats.items())))
print(f'Tier2 V4 cluster {cluster.key}: ' + ' '.join(f'{k}={v}' for k, v in sorted(stats.items())))
total.update(stats)
hooks_by_unit: Dict[int, List[Tuple[Cluster, int]]] = collections.defaultdict(list)
@@ -802,7 +914,7 @@ def main() -> int:
# common file is handwritten by V2 and must remain; generated cluster files
# carry all hot code.
print('Tier2 SUPERBLOCK V3 DATAFLOW:',
print('Tier2 SUPERBLOCK V4 150FPS:',
f'clusters={len(CLUSTERS)} hook_units_changed={hook_units_changed}',
f'blocks={total["blocks"]} lines={total["lines"]}',
f'fused_calls={total["fused_calls"]} fused_tail={total["fused_tail"]}',