From 1d348ac2d37eda5380165315a40216de3f4e78cb Mon Sep 17 00:00:00 2001 From: Will Martin Date: Thu, 29 Jan 2026 22:28:05 +0900 Subject: [PATCH] [Metal] Optimize transfer path and RT handling - Add tile-instanced transfer batching and per-frame instance buffers - Specialize transfer shaders and gate sample-id usage for MSAA - Add blit fast paths for compatible color/depth transfers - Fix dummy RT MSAA keying and avoid unnecessary load/store - Fix stencil view lifetime and depth+stencil transfer packing - Reduce resolve/readback stalls and refresh render-pass/scissor state --- src/xenia/gpu/metal/dxbc_to_dxil_converter.cc | 6 +- src/xenia/gpu/metal/dxbc_to_dxil_converter.h | 9 + src/xenia/gpu/metal/ir_runtime_impl.mm | 2 +- .../gpu/metal/metal_command_processor.cc | 1039 +++++++++-- src/xenia/gpu/metal/metal_command_processor.h | 70 +- src/xenia/gpu/metal/metal_geometry_shader.cc | 2 +- src/xenia/gpu/metal/metal_geometry_shader.h | 2 +- src/xenia/gpu/metal/metal_graphics_system.cc | 2 +- src/xenia/gpu/metal/metal_graphics_system.h | 2 +- src/xenia/gpu/metal/metal_heap_pool.cc | 2 +- src/xenia/gpu/metal/metal_heap_pool.h | 2 +- .../gpu/metal/metal_primitive_processor.cc | 2 +- .../gpu/metal/metal_primitive_processor.h | 2 +- .../gpu/metal/metal_render_target_cache.cc | 1616 ++++++++++++++--- .../gpu/metal/metal_render_target_cache.h | 43 +- src/xenia/gpu/metal/metal_shader.cc | 2 +- src/xenia/gpu/metal/metal_shader.h | 2 +- src/xenia/gpu/metal/metal_shader_cache.cc | 7 +- src/xenia/gpu/metal/metal_shader_cache.h | 7 +- src/xenia/gpu/metal/metal_shader_converter.cc | 2 +- src/xenia/gpu/metal/metal_shader_converter.h | 2 +- src/xenia/gpu/metal/metal_shared_memory.cc | 9 + src/xenia/gpu/metal/metal_shared_memory.h | 2 +- src/xenia/gpu/metal/metal_texture_cache.cc | 82 +- src/xenia/gpu/metal/metal_texture_cache.h | 16 +- src/xenia/gpu/metal/metal_trace_dump_main.cc | 2 +- .../gpu/metal/metal_trace_viewer_main.cc | 2 +- 27 files changed, 2413 insertions(+), 523 deletions(-) diff --git a/src/xenia/gpu/metal/dxbc_to_dxil_converter.cc b/src/xenia/gpu/metal/dxbc_to_dxil_converter.cc index 9d9a813a0..0bc3e921e 100644 --- a/src/xenia/gpu/metal/dxbc_to_dxil_converter.cc +++ b/src/xenia/gpu/metal/dxbc_to_dxil_converter.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -88,9 +88,9 @@ bool DxbcToDxilConverter::Initialize() { dxilconv_path_ = "linked"; is_available_ = true; if (extra_options && *extra_options) { - XELOGI("DxbcToDxilConverter: Using extra options: {}", extra_options); + XELOGD("DxbcToDxilConverter: Using extra options: {}", extra_options); } else if (extra_options && !*extra_options) { - XELOGI("DxbcToDxilConverter: Extra options disabled via env"); + XELOGD("DxbcToDxilConverter: Extra options disabled via env"); } else { XELOGI( "DxbcToDxilConverter: Using default extra options: " diff --git a/src/xenia/gpu/metal/dxbc_to_dxil_converter.h b/src/xenia/gpu/metal/dxbc_to_dxil_converter.h index 526ae7c93..24c08c446 100644 --- a/src/xenia/gpu/metal/dxbc_to_dxil_converter.h +++ b/src/xenia/gpu/metal/dxbc_to_dxil_converter.h @@ -1,3 +1,12 @@ +/** + ****************************************************************************** + * Xenia : Xbox 360 Emulator Research Project * + ****************************************************************************** + * Copyright 2026 Ben Vanik. All rights reserved. * + * Released under the BSD license - see LICENSE in the root for more details. * + ****************************************************************************** + */ + /** * DXBC to DXIL converter wrapper for Metal backend. * diff --git a/src/xenia/gpu/metal/ir_runtime_impl.mm b/src/xenia/gpu/metal/ir_runtime_impl.mm index d743808e1..c83cb1b32 100644 --- a/src/xenia/gpu/metal/ir_runtime_impl.mm +++ b/src/xenia/gpu/metal/ir_runtime_impl.mm @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_command_processor.cc b/src/xenia/gpu/metal/metal_command_processor.cc index 4a3686564..80f045b8d 100644 --- a/src/xenia/gpu/metal/metal_command_processor.cc +++ b/src/xenia/gpu/metal/metal_command_processor.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -16,6 +16,7 @@ #include #include #include +#include #include #include #include @@ -37,6 +38,7 @@ #include "xenia/base/filesystem.h" #include "xenia/base/logging.h" #include "xenia/base/math.h" +#include "xenia/base/memory.h" #include "xenia/base/profiling.h" #include "xenia/base/xxhash.h" #include "xenia/gpu/draw_util.h" @@ -47,9 +49,11 @@ #include "xenia/gpu/metal/metal_shader_converter.h" #include "xenia/gpu/packet_disassembler.h" #include "xenia/gpu/registers.h" +#include "xenia/gpu/texture_info.h" #include "xenia/gpu/xenos.h" #include "xenia/kernel/kernel_state.h" #include "xenia/kernel/user_module.h" +#include "xenia/ui/metal/metal_gpu_completion_timeline.h" using BYTE = uint8_t; #include "xenia/gpu/shaders/bytecode/d3d12_5_1/adaptive_quad_hs.h" #include "xenia/gpu/shaders/bytecode/d3d12_5_1/adaptive_triangle_hs.h" @@ -63,6 +67,7 @@ using BYTE = uint8_t; #include "xenia/gpu/shaders/bytecode/d3d12_5_1/discrete_triangle_3cp_hs.h" #include "xenia/gpu/shaders/bytecode/d3d12_5_1/tessellation_adaptive_vs.h" #include "xenia/gpu/shaders/bytecode/d3d12_5_1/tessellation_indexed_vs.h" +#include "xenia/gpu/shaders/bytecode/metal/resolve_downscale_cs.h" #include "xenia/ui/metal/metal_presenter.h" #ifndef DISPATCH_DATA_DESTRUCTOR_NONE @@ -77,6 +82,7 @@ using BYTE = uint8_t; // kIRDescriptorHeapBindPoint = 0 DECLARE_bool(clear_memory_page_state); +DECLARE_bool(submit_on_primary_buffer_end); // kIRSamplerHeapBindPoint = 1 // kIRArgumentBufferBindPoint = 2 // kIRArgumentBufferDrawArgumentsBindPoint = 4 @@ -116,6 +122,47 @@ void LogMetalErrorDetails(const char* label, NS::Error* error) { } } +MTL::ComputePipelineState* CreateComputePipelineFromEmbeddedLibrary( + MTL::Device* device, const void* metallib_data, size_t metallib_size, + const char* debug_name) { + if (!device || !metallib_data || !metallib_size) { + return nullptr; + } + + NS::Error* error = nullptr; + dispatch_data_t data = dispatch_data_create( + metallib_data, metallib_size, nullptr, DISPATCH_DATA_DESTRUCTOR_DEFAULT); + MTL::Library* lib = device->newLibrary(data, &error); + dispatch_release(data); + if (!lib) { + XELOGE("Metal: failed to create {} library: {}", debug_name, + error ? error->localizedDescription()->utf8String() : "unknown"); + return nullptr; + } + + // XeSL compute entrypoint name used in the embedded metallibs. + NS::String* fn_name = NS::String::string("entry_xe", NS::UTF8StringEncoding); + MTL::Function* fn = lib->newFunction(fn_name); + if (!fn) { + XELOGE("Metal: {} missing entry_xe", debug_name); + lib->release(); + return nullptr; + } + + MTL::ComputePipelineState* pipeline = + device->newComputePipelineState(fn, &error); + fn->release(); + lib->release(); + + if (!pipeline) { + XELOGE("Metal: failed to create {} pipeline: {}", debug_name, + error ? error->localizedDescription()->utf8String() : "unknown"); + return nullptr; + } + + return pipeline; +} + constexpr uint32_t kPipelineDiskCacheMagic = 0x43504D58; // 'XMPC' constexpr uint32_t kPipelineDiskCacheVersion = 2; constexpr size_t kPipelineDiskCacheMaxEntrySize = 1 << 20; @@ -273,6 +320,61 @@ MTL::BlendFactor ToMetalBlendFactorAlpha(xenos::BlendFactor blend_factor) { return kBlendFactorAlphaMap[uint32_t(blend_factor) & 0x1F]; } +void DownscaleResolveTileData(const uint8_t* source, uint8_t* dest, + uint32_t tile_count, uint32_t pixel_size_log2, + uint32_t scale_x, uint32_t scale_y, + bool half_pixel_offset) { + if (!source || !dest || tile_count == 0) { + return; + } + const uint32_t pixel_size = 1u << pixel_size_log2; + const uint32_t scale_xy = scale_x * scale_y; + const uint32_t tile_size_1x = 32u * 32u * pixel_size; + const uint32_t tile_size_scaled = tile_size_1x * scale_xy; + const uint32_t block_sample_offset = + (half_pixel_offset && scale_xy > 1u) + ? ((scale_x >> 1u) + (scale_y >> 1u) * scale_x) + : 0u; + const uint32_t block_sample_offset_bytes = block_sample_offset * pixel_size; + const uint32_t src_pixel_stride = pixel_size * scale_xy; + + for (uint32_t tile_index = 0; tile_index < tile_count; ++tile_index) { + const uint8_t* src_tile = + source + tile_index * tile_size_scaled + block_sample_offset_bytes; + uint8_t* dst_tile = dest + tile_index * tile_size_1x; + for (uint32_t pixel_index = 0; pixel_index < 32u * 32u; ++pixel_index) { + const uint32_t src_offset = pixel_index * src_pixel_stride; + const uint32_t dst_offset = pixel_index * pixel_size; + switch (pixel_size_log2) { + case 0: { + dst_tile[dst_offset] = src_tile[src_offset]; + break; + } + case 1: { + uint16_t value; + std::memcpy(&value, src_tile + src_offset, sizeof(value)); + std::memcpy(dst_tile + dst_offset, &value, sizeof(value)); + break; + } + case 2: { + uint32_t value; + std::memcpy(&value, src_tile + src_offset, sizeof(value)); + std::memcpy(dst_tile + dst_offset, &value, sizeof(value)); + break; + } + case 3: { + uint64_t value; + std::memcpy(&value, src_tile + src_offset, sizeof(value)); + std::memcpy(dst_tile + dst_offset, &value, sizeof(value)); + break; + } + default: + break; + } + } + } +} + } // namespace MetalCommandProcessor::MetalCommandProcessor( @@ -401,6 +503,115 @@ MetalCommandProcessor::DrawRingBuffers::~DrawRingBuffers() { } } +void MetalCommandProcessor::UpdateDebugMarkersEnabled() { + // Enable debug markers if the CVAR is set (RenderDoc auto-detect disabled on + // macOS). + debug_markers_enabled_ = IsGpuDebugMarkersEnabled(); +} + +void MetalCommandProcessor::PushDebugMarker(const char* format, ...) { + if (!debug_markers_enabled_) { + return; + } + char label[256]; + va_list args; + va_start(args, format); + vsnprintf(label, sizeof(label), format, args); + va_end(args); + auto* ns_label = NS::String::string(label, NS::UTF8StringEncoding); + if (current_render_encoder_) { + current_render_encoder_->pushDebugGroup(ns_label); + debug_marker_stack_.push_back(DebugMarkerTarget::kRenderEncoder); + } else if (current_command_buffer_) { + current_command_buffer_->pushDebugGroup(ns_label); + debug_marker_stack_.push_back(DebugMarkerTarget::kCommandBuffer); + } +} + +void MetalCommandProcessor::PopDebugMarker() { + if (!debug_markers_enabled_ || debug_marker_stack_.empty()) { + return; + } + DebugMarkerTarget target = debug_marker_stack_.back(); + debug_marker_stack_.pop_back(); + if (target == DebugMarkerTarget::kRenderEncoder) { + if (current_render_encoder_) { + current_render_encoder_->popDebugGroup(); + } + } else { + if (current_command_buffer_) { + current_command_buffer_->popDebugGroup(); + } + } +} + +void MetalCommandProcessor::InsertDebugMarker(const char* format, ...) { + if (!debug_markers_enabled_) { + return; + } + char label[256]; + va_list args; + va_start(args, format); + vsnprintf(label, sizeof(label), format, args); + va_end(args); + auto* ns_label = NS::String::string(label, NS::UTF8StringEncoding); + if (current_render_encoder_) { + current_render_encoder_->insertDebugSignpost(ns_label); + } else if (current_command_buffer_) { + current_command_buffer_->pushDebugGroup(ns_label); + current_command_buffer_->popDebugGroup(); + } +} + +void MetalCommandProcessor::RequestCapture() { + capture_requested_.store(true, std::memory_order_release); +} + +void MetalCommandProcessor::MaybeStartCapture() { + if (!capture_requested_.exchange(false, std::memory_order_acq_rel)) { + return; + } + if (!command_queue_) { + XELOGW("Metal capture requested but command queue is not ready"); + return; + } + capture_manager_ = MTL::CaptureManager::sharedCaptureManager(); + if (!capture_manager_) { + XELOGW("Metal capture manager not available"); + return; + } + auto* descriptor = MTL::CaptureDescriptor::alloc()->init(); + descriptor->setCaptureObject(command_queue_); + descriptor->setDestination(MTL::CaptureDestinationGPUTraceDocument); + + const char* capture_dir = std::getenv("XENIA_GPU_CAPTURE_DIR"); + std::string filename = + capture_dir ? (std::string(capture_dir) + "/metal_capture.gputrace") + : std::string("./metal_capture.gputrace"); + auto* url = NS::URL::fileURLWithPath( + NS::String::string(filename.c_str(), NS::UTF8StringEncoding)); + descriptor->setOutputURL(url); + + NS::Error* error = nullptr; + if (capture_manager_->startCapture(descriptor, &error)) { + capture_active_ = true; + XELOGI("Metal capture started: {}", filename); + } else { + XELOGE("Metal capture start failed: {} (set MTL_CAPTURE_ENABLED=1)", + error ? error->localizedDescription()->utf8String() : "unknown"); + } + descriptor->release(); +} + +void MetalCommandProcessor::StopCaptureIfActive() { + if (!capture_active_ || !capture_manager_) { + return; + } + capture_manager_->stopCapture(); + capture_active_ = false; + XELOGI("Metal capture completed"); +} + void MetalCommandProcessor::TracePlaybackWroteMemory(uint32_t base_ptr, uint32_t length) { if (shared_memory_) { @@ -445,14 +656,58 @@ void MetalCommandProcessor::InvalidateGpuMemory() { } void MetalCommandProcessor::ClearReadbackBuffers() { - // TODO(wmarti): Implement readback buffer clearing when memexport readback - // is added. See D3D12's readback_buffers_ and memexport_readback_buffers_. + for (auto& entry : readback_buffers_) { + ReadbackBuffer& rb = entry.second; + for (size_t i = 0; i < 2; ++i) { + if (rb.buffers[i]) { + rb.buffers[i]->release(); + rb.buffers[i] = nullptr; + } + rb.sizes[i] = 0; + rb.submission_ids[i] = 0; + } + rb.current_index = 0; + rb.last_used_frame = 0; + } + readback_buffers_.clear(); +} + +void MetalCommandProcessor::EvictOldReadbackBuffers( + std::unordered_map& buffer_map) { + if (frame_current_ <= kReadbackBufferEvictionAgeFrames) { + return; + } + + for (auto it = buffer_map.begin(); it != buffer_map.end();) { + if (it->second.last_used_frame < + frame_current_ - kReadbackBufferEvictionAgeFrames) { + for (int i = 0; i < 2; ++i) { + if (it->second.buffers[i]) { + it->second.buffers[i]->release(); + } + } + it = buffer_map.erase(it); + } else { + ++it; + } + } } ui::metal::MetalProvider& MetalCommandProcessor::GetMetalProvider() const { return *static_cast(graphics_system_->provider()); } +uint64_t MetalCommandProcessor::GetCurrentSubmission() const { + return completion_timeline_ ? completion_timeline_->GetUpcomingSubmission() + : 1; +} + +uint64_t MetalCommandProcessor::GetCompletedSubmission() const { + return completion_timeline_ + ? completion_timeline_->GetCompletedSubmissionFromLastUpdate() + : 0; +} + void MetalCommandProcessor::MarkResolvedMemory(uint32_t base_ptr, uint32_t length) { if (length == 0) return; @@ -522,6 +777,12 @@ bool MetalCommandProcessor::SetupContext() { return false; } + // Check if debug markers should be enabled (CVAR). + UpdateDebugMarkersEnabled(); + if (debug_markers_enabled_) { + XELOGI("GPU debug markers enabled for Metal debug tools"); + } + const ui::metal::MetalProvider& provider = GetMetalProvider(); device_ = provider.GetDevice(); command_queue_ = provider.GetCommandQueue(); @@ -542,6 +803,18 @@ bool MetalCommandProcessor::SetupContext() { "waitUntilCompleted"); } + completion_timeline_ = ui::metal::MetalGPUCompletionTimeline::Create(device_); + if (!completion_timeline_) { + XELOGE("MetalCommandProcessor: Failed to create completion timeline"); + return false; + } + submission_open_ = false; + submission_completed_processed_ = 0; + frame_open_ = false; + frame_current_ = 1; + frame_completed_ = 0; + std::fill_n(closed_frame_submissions_, kQueueFrames, 0); + bool supports_apple7 = device_->supportsFamily(MTL::GPUFamilyApple7); bool supports_mac2 = device_->supportsFamily(MTL::GPUFamilyMac2); mesh_shader_supported_ = supports_apple7 || supports_mac2; @@ -563,8 +836,26 @@ bool MetalCommandProcessor::SetupContext() { return false; } - texture_cache_ = std::make_unique(this, *register_file_, - *shared_memory_, 1, 1); + // Get the draw resolution scale for the render target cache and the texture + // cache (match D3D12/Vulkan behavior). + uint32_t draw_resolution_scale_x = 1; + uint32_t draw_resolution_scale_y = 1; + bool draw_resolution_scale_not_clamped = + TextureCache::GetConfigDrawResolutionScale(draw_resolution_scale_x, + draw_resolution_scale_y); + if (!draw_resolution_scale_not_clamped) { + XELOGW( + "The requested draw resolution scale is not supported by the emulator " + "or config, reducing to {}x{}", + draw_resolution_scale_x, draw_resolution_scale_y); + } + XELOGI("Metal: draw resolution scale {}x{} (supported={})", + draw_resolution_scale_x, draw_resolution_scale_y, + draw_resolution_scale_not_clamped); + + texture_cache_ = std::make_unique( + this, *register_file_, *shared_memory_, draw_resolution_scale_x, + draw_resolution_scale_y); if (!texture_cache_->Initialize()) { XELOGE("Failed to initialize Metal texture cache"); return false; @@ -572,12 +863,20 @@ bool MetalCommandProcessor::SetupContext() { // Initialize render target cache render_target_cache_ = std::make_unique( - *register_file_, *memory_, &trace_writer_, 1, 1, *this); + *register_file_, *memory_, &trace_writer_, draw_resolution_scale_x, + draw_resolution_scale_y, *this); if (!render_target_cache_->Initialize()) { XELOGE("Failed to initialize Metal render target cache"); return false; } + resolve_downscale_pipeline_ = CreateComputePipelineFromEmbeddedLibrary( + device_, resolve_downscale_cs_metallib, + sizeof(resolve_downscale_cs_metallib), "resolve_downscale"); + if (!resolve_downscale_pipeline_) { + XELOGW("MetalCommandProcessor: resolve downscale pipeline unavailable"); + } + // Initialize shader translation pipeline if (!InitializeShaderTranslation()) { XELOGE("Failed to initialize shader translation"); @@ -814,31 +1113,21 @@ void MetalCommandProcessor::PrepareForWait() { // By submitting and waiting for all GPU work now, we ensure clean pool // drainage. - if (current_render_encoder_) { - current_render_encoder_->endEncoding(); - current_render_encoder_->release(); - current_render_encoder_ = nullptr; - } + EndRenderEncoder(); - if (current_command_buffer_) { - uint64_t wait_value = 0; - if (wait_shared_event_) { - wait_value = ++wait_shared_event_value_; - current_command_buffer_->encodeSignalEvent(wait_shared_event_, - wait_value); + if (submission_open_ || current_command_buffer_) { + uint64_t submission_to_wait = + current_command_buffer_ ? GetCurrentSubmission() : 0; + if (!submission_open_) { + XELOGW( + "MetalCommandProcessor::PrepareForWait: command buffer without " + "open submission"); + submission_open_ = true; } - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - if (wait_shared_event_) { - wait_shared_event_->waitUntilSignaledValue( - wait_value, std::numeric_limits::max()); - } else { - current_command_buffer_->waitUntilCompleted(); + EndSubmission(false); + if (submission_to_wait) { + CheckSubmissionCompletion(submission_to_wait); } - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; } DrainCommandBufferAutoreleasePool(); @@ -875,32 +1164,21 @@ void MetalCommandProcessor::PrepareForWait() { } void MetalCommandProcessor::ShutdownContext() { - // End any active render encoder before shutdown - if (current_render_encoder_) { - current_render_encoder_->endEncoding(); - // Don't release yet - wait until command buffer completes - } + EndRenderEncoder(); - // Submit and wait for any pending command buffer - if (current_command_buffer_) { - uint64_t wait_value = 0; - if (wait_shared_event_) { - wait_value = ++wait_shared_event_value_; - current_command_buffer_->encodeSignalEvent(wait_shared_event_, - wait_value); + if (submission_open_ || current_command_buffer_) { + uint64_t submission_to_wait = + current_command_buffer_ ? GetCurrentSubmission() : 0; + if (!submission_open_) { + XELOGW( + "MetalCommandProcessor::ShutdownContext: command buffer without " + "open submission"); + submission_open_ = true; } - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - if (wait_shared_event_) { - wait_shared_event_->waitUntilSignaledValue( - wait_value, std::numeric_limits::max()); - } else { - current_command_buffer_->waitUntilCompleted(); + EndSubmission(false); + if (submission_to_wait) { + CheckSubmissionCompletion(submission_to_wait); } - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; } // Even if we have no active command buffer at this point, there may be @@ -938,6 +1216,17 @@ void MetalCommandProcessor::ShutdownContext() { } DrainCommandBufferAutoreleasePool(); + ClearReadbackBuffers(); + if (resolve_downscale_buffer_) { + resolve_downscale_buffer_->release(); + resolve_downscale_buffer_ = nullptr; + resolve_downscale_buffer_size_ = 0; + } + if (resolve_downscale_pipeline_) { + resolve_downscale_pipeline_->release(); + resolve_downscale_pipeline_ = nullptr; + } + { std::lock_guard lock(draw_ring_mutex_); active_draw_ring_.reset(); @@ -964,6 +1253,12 @@ void MetalCommandProcessor::ShutdownContext() { } depth_only_pixel_function_name_.clear(); frame_open_ = false; + frame_current_ = 1; + frame_completed_ = 0; + std::fill_n(closed_frame_submissions_, kQueueFrames, 0); + submission_open_ = false; + submission_completed_processed_ = 0; + completion_timeline_.reset(); shader_cache_.clear(); shared_memory_.reset(); @@ -1447,27 +1742,7 @@ void MetalCommandProcessor::IssueSwap(uint32_t frontbuffer_ptr, last_swap_ptr_ = frontbuffer_ptr; last_swap_width_ = frontbuffer_width; last_swap_height_ = frontbuffer_height; - - // End any active render encoder - EndRenderEncoder(); - - // Submit and wait for command buffer - if (current_command_buffer_) { - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; - } - - if (primitive_processor_ && frame_open_) { - primitive_processor_->EndFrame(); - frame_open_ = false; - } - if (shared_memory_ && ::cvars::clear_memory_page_state) { - shared_memory_->SetSystemPageBlocksValidWithGpuDataWritten(); - } + EndSubmission(true); // Push the rendered frame to the presenter's guest output mailbox // This is required for trace dumps to capture the output. Use the @@ -1566,8 +1841,9 @@ void MetalCommandProcessor::IssueSwap(uint32_t frontbuffer_ptr, uint32_t source_height = output_height; bool force_swap_rb_copy = force_swap_rb; bool use_pwl_gamma_ramp_copy = use_pwl_gamma_ramp; + auto aspect = graphics_system_->GetScaledAspectRatio(); presenter->RefreshGuestOutput( - output_width, output_height, 1280, 720, // Display aspect ratio + output_width, output_height, aspect.first, aspect.second, [source_texture, metal_presenter, source_width, source_height, force_swap_rb_copy, use_pwl_gamma_ramp_copy]( ui::Presenter::GuestOutputRefreshContext& context) -> bool { @@ -1575,13 +1851,27 @@ void MetalCommandProcessor::IssueSwap(uint32_t frontbuffer_ptr, static_cast( context); context.SetIs8bpc(!use_pwl_gamma_ramp_copy); - return metal_presenter->CopyTextureToGuestOutput( + uint64_t submission_id = 0; + bool copy_ok = metal_presenter->CopyTextureToGuestOutput( source_texture, metal_context.resource_uav_capable(), source_width, source_height, force_swap_rb_copy, - use_pwl_gamma_ramp_copy); + use_pwl_gamma_ramp_copy, &submission_id); + if (submission_id) { + metal_context.SetSubmissionId(submission_id); + } + return copy_ok; }); } } + + StopCaptureIfActive(); +} + +void MetalCommandProcessor::OnPrimaryBufferEnd() { + if (cvars::submit_on_primary_buffer_end && submission_open_ && + CanEndSubmissionImmediately()) { + EndSubmission(false); + } } Shader* MetalCommandProcessor::LoadShader(xenos::ShaderType shader_type, @@ -1619,6 +1909,10 @@ bool MetalCommandProcessor::IssueDraw(xenos::PrimitiveType primitive_type, const RegisterFile& regs = *register_file_; uint32_t normalized_color_mask = 0; + if (!BeginSubmission(true)) { + return false; + } + // Check for copy mode xenos::EdramMode edram_mode = regs.Get().edram_mode; if (edram_mode == xenos::EdramMode::kCopy) { @@ -2889,6 +3183,10 @@ bool MetalCommandProcessor::IssueDraw(xenos::PrimitiveType primitive_type, } bool MetalCommandProcessor::IssueCopy() { + if (!BeginSubmission(true)) { + return false; + } + // Finish any in-flight rendering so the render target contents are // available to the render target cache, similar to D3D12's // D3D12CommandProcessor::IssueCopy. @@ -2899,18 +3197,10 @@ bool MetalCommandProcessor::IssueCopy() { } if (!current_command_buffer_) { - if (!command_queue_) { - XELOGE("MetalCommandProcessor::IssueCopy: no command queue"); + if (!EnsureCommandBuffer()) { + XELOGE("MetalCommandProcessor::IssueCopy: no command buffer"); return false; } - // Note: commandBuffer() returns an autoreleased object, we must retain it. - current_command_buffer_ = command_queue_->commandBuffer(); - if (!current_command_buffer_) { - XELOGE( - "MetalCommandProcessor::IssueCopy: failed to create command buffer"); - return false; - } - current_command_buffer_->retain(); current_command_buffer_->setLabel( NS::String::string("XeniaCopyCommandBuffer", NS::UTF8StringEncoding)); } @@ -2929,14 +3219,57 @@ bool MetalCommandProcessor::IssueCopy() { return false; } + ReadbackResolveMode readback_mode = GetReadbackResolveMode(); + bool do_readback = (readback_mode != ReadbackResolveMode::kDisabled); + bool readback_scaled = false; + bool readback_scaled_gpu = false; + bool use_gpu_downscale = false; + bool readback_scheduled = false; + ReadbackBuffer* readback_buffer = nullptr; + uint32_t write_index = 0; + uint32_t read_index = 0; + bool use_delayed_sync = false; + bool wait_for_completion = false; + bool should_copy = false; + bool is_cache_miss = false; + uint32_t source_length = 0; + uint32_t readback_length = 0; + uint32_t tile_count = 0; + uint32_t pixel_size_log2 = 0; + uint32_t scale_x = 1; + uint32_t scale_y = 1; + bool half_pixel_offset = false; + uint32_t source_offset_bytes = 0; + uint64_t scaled_range_offset_bytes = 0; + uint64_t readback_base_offset_bytes = 0; + uint64_t scaled_copy_length = 0; + size_t source_buffer_binding_offset = 0; + uint64_t source_offset_bytes_log = 0; + + if (do_readback) { + // Early check: if destination memory is not accessible, skip readback. + VirtualHeap* physical_heap = memory_->GetPhysicalHeap(); + bool memory_accessible = false; + if (physical_heap) { + HeapAllocationInfo alloc_info; + if (physical_heap->QueryRegionInfo(written_address, &alloc_info) && + (alloc_info.state & kMemoryAllocationCommit) && + (alloc_info.protect & kMemoryProtectWrite)) { + uint32_t end_address = written_address + written_length; + uint32_t region_end = alloc_info.base_address + alloc_info.region_size; + if (end_address <= region_end) { + memory_accessible = true; + } + } + } + if (!memory_accessible) { + do_readback = false; + } + } + if (!written_length) { // Commit any in-flight work so ordering matches D3D12 submission behavior. - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; + EndSubmission(false); return true; } @@ -2944,27 +3277,320 @@ bool MetalCommandProcessor::IssueCopy() { // with stale MemoryRead commands from the trace file. MarkResolvedMemory(written_address, written_length); - // The resolve writes into guest memory via the shared memory buffer. - // Any cached views of this memory (especially textures sourced from it) - // must be invalidated, otherwise subsequent render-to-texture / postprocess - // passes will sample stale host textures and produce corrupted output. - // if (shared_memory_) { - // shared_memory_->MemoryInvalidationCallback(written_address, - // written_length, true); - // } - // if (primitive_processor_) { - // primitive_processor_->MemoryInvalidationCallback(written_address, - // written_length, true); - // } + if (do_readback) { + MTL::Buffer* source_buffer = nullptr; + size_t source_offset = 0; + size_t source_length_size_t = 0; + + if (texture_cache_ && texture_cache_->IsDrawResolutionScaled()) { + readback_scaled = true; + auto* metal_texture_cache = + static_cast(texture_cache_.get()); + if (!metal_texture_cache || + !metal_texture_cache->GetCurrentScaledResolveBuffer( + source_buffer, source_offset, source_length_size_t)) { + XELOGE("MetalResolveReadback: failed to get scaled resolve buffer"); + do_readback = false; + } else { + scale_x = texture_cache_->draw_resolution_scale_x(); + scale_y = texture_cache_->draw_resolution_scale_y(); + uint64_t scale_area = uint64_t(scale_x) * uint64_t(scale_y); + uint64_t range_start_scaled = + metal_texture_cache->GetCurrentScaledResolveRangeStartScaled(); + if (scale_area && (range_start_scaled % scale_area) == 0) { + uint64_t range_start_unscaled = range_start_scaled / scale_area; + if (written_address >= range_start_unscaled) { + scaled_range_offset_bytes = + (uint64_t(written_address) - range_start_unscaled) * scale_area; + } + } + if (scaled_range_offset_bytes > source_length_size_t) { + XELOGE("MetalResolveReadback: scaled range offset out of bounds"); + do_readback = false; + } else { + readback_base_offset_bytes = scaled_range_offset_bytes; + } + } + } else { + source_buffer = shared_memory_ ? shared_memory_->GetBuffer() : nullptr; + source_offset = written_address; + source_length_size_t = written_length; + source_offset_bytes_log = source_offset; + } + + if (do_readback) { + if (!source_buffer || source_length_size_t == 0) { + do_readback = false; + } else if (source_length_size_t > std::numeric_limits::max()) { + XELOGE("MetalResolveReadback: source length too large ({})", + source_length_size_t); + do_readback = false; + } else if (source_offset + source_length_size_t > + size_t(source_buffer->length())) { + XELOGE("MetalResolveReadback: source range out of bounds"); + do_readback = false; + } + } + + if (do_readback) { + source_length = uint32_t(source_length_size_t); + uint64_t scaled_available_bytes = source_length_size_t; + if (readback_scaled) { + if (scaled_range_offset_bytes >= source_length_size_t) { + XELOGE("MetalResolveReadback: scaled range offset exceeds length"); + do_readback = false; + } else { + scaled_available_bytes = + source_length_size_t - scaled_range_offset_bytes; + } + } + uint64_t resolve_key = + MakeReadbackResolveKey(written_address, written_length); + ReadbackBuffer& rb = readback_buffers_[resolve_key]; + rb.last_used_frame = frame_current_; + readback_buffer = &rb; + write_index = rb.current_index; + use_delayed_sync = (readback_mode == ReadbackResolveMode::kFast || + readback_mode == ReadbackResolveMode::kSome); + read_index = use_delayed_sync ? (1u - write_index) : write_index; + + readback_length = source_length; + if (readback_scaled) { + auto copy_dest_info = register_file_->Get(); + const FormatInfo* format_info = + FormatInfo::Get(uint32_t(copy_dest_info.copy_dest_format)); + uint32_t bits_per_pixel = format_info->bits_per_pixel; + xe::bit_scan_forward(bits_per_pixel >> 3, &pixel_size_log2); + uint32_t bytes_per_pixel = 1u << pixel_size_log2; + uint32_t tile_size_1x = 32u * 32u * bytes_per_pixel; + tile_count = written_length / tile_size_1x; + half_pixel_offset = cvars::readback_resolve_half_pixel_offset && + (scale_x > 1 || scale_y > 1); + uint64_t tile_size_scaled = + uint64_t(tile_size_1x) * uint64_t(scale_x) * uint64_t(scale_y); + uint64_t required_scaled = uint64_t(tile_count) * tile_size_scaled; + if (required_scaled > scaled_available_bytes) { + tile_count = uint32_t(scaled_available_bytes / tile_size_scaled); + required_scaled = uint64_t(tile_count) * tile_size_scaled; + } + if (tile_count == 0) { + do_readback = false; + } + + uint64_t source_offset_bytes_64 = + uint64_t(source_offset) + scaled_range_offset_bytes; + source_offset_bytes_log = source_offset_bytes_64; + if (do_readback && tile_count != 0 && resolve_downscale_pipeline_) { + uint32_t downscale_buffer_size = + AlignReadbackBufferSize(written_length); + if (downscale_buffer_size > resolve_downscale_buffer_size_) { + if (resolve_downscale_buffer_) { + resolve_downscale_buffer_->release(); + resolve_downscale_buffer_ = nullptr; + resolve_downscale_buffer_size_ = 0; + } + if (device_) { + resolve_downscale_buffer_ = device_->newBuffer( + downscale_buffer_size, MTL::ResourceStorageModePrivate); + if (resolve_downscale_buffer_) { + resolve_downscale_buffer_size_ = downscale_buffer_size; + } + } + } + if (resolve_downscale_buffer_) { + use_gpu_downscale = true; + readback_length = written_length; + source_buffer_binding_offset = + size_t(source_offset_bytes_64 & ~uint64_t(3)); + source_offset_bytes = + uint32_t(source_offset_bytes_64 - + uint64_t(source_buffer_binding_offset)); + } + } + if (do_readback && !use_gpu_downscale) { + scaled_copy_length = required_scaled; + if (scaled_copy_length > std::numeric_limits::max()) { + XELOGE("MetalResolveReadback: scaled copy length too large ({})", + scaled_copy_length); + do_readback = false; + } else { + readback_length = uint32_t(scaled_copy_length); + source_offset = size_t(source_offset_bytes_64); + readback_base_offset_bytes = 0; + } + } + } + + uint32_t aligned_size = AlignReadbackBufferSize(readback_length); + if (aligned_size > rb.sizes[write_index]) { + if (!device_) { + XELOGE("MetalResolveReadback: missing Metal device"); + do_readback = false; + } + } + if (do_readback && aligned_size > rb.sizes[write_index]) { + if (rb.buffers[write_index]) { + rb.buffers[write_index]->release(); + rb.buffers[write_index] = nullptr; + } + rb.buffers[write_index] = + device_->newBuffer(aligned_size, MTL::ResourceStorageModeShared); + rb.sizes[write_index] = aligned_size; + rb.submission_ids[write_index] = 0; + } + if (do_readback && !rb.buffers[write_index]) { + XELOGE("MetalResolveReadback: failed to allocate readback buffer"); + do_readback = false; + } else if (do_readback) { + if (readback_scaled && use_gpu_downscale) { + MTL::ComputeCommandEncoder* compute = + current_command_buffer_->computeCommandEncoder(); + if (!compute) { + XELOGE("MetalResolveReadback: failed to create compute encoder"); + do_readback = false; + } else { + struct ResolveDownscaleConstants { + uint32_t scale_x; + uint32_t scale_y; + uint32_t pixel_size_log2; + uint32_t tile_count; + uint32_t source_offset_bytes; + uint32_t half_pixel_offset; + } constants; + constants.scale_x = scale_x; + constants.scale_y = scale_y; + constants.pixel_size_log2 = pixel_size_log2; + constants.tile_count = tile_count; + constants.source_offset_bytes = source_offset_bytes; + constants.half_pixel_offset = half_pixel_offset ? 1u : 0u; + + compute->setComputePipelineState(resolve_downscale_pipeline_); + compute->setBytes(&constants, sizeof(constants), 0); + compute->setBuffer(source_buffer, source_buffer_binding_offset, 1); + compute->setBuffer(resolve_downscale_buffer_, 0, 2); + compute->useResource(source_buffer, MTL::ResourceUsageRead); + compute->useResource(resolve_downscale_buffer_, + MTL::ResourceUsageWrite); + compute->dispatchThreadgroups(MTL::Size::Make(tile_count, 1, 1), + MTL::Size::Make(32, 32, 1)); + compute->endEncoding(); + + MTL::BlitCommandEncoder* blit = + current_command_buffer_->blitCommandEncoder(); + if (!blit) { + XELOGE("MetalResolveReadback: failed to create blit encoder"); + do_readback = false; + } else { + blit->copyFromBuffer(resolve_downscale_buffer_, 0, + rb.buffers[write_index], 0, readback_length); + blit->endEncoding(); + rb.submission_ids[write_index] = GetCurrentSubmission(); + readback_scheduled = true; + readback_scaled_gpu = true; + } + } + } else { + MTL::BlitCommandEncoder* blit = + current_command_buffer_->blitCommandEncoder(); + if (!blit) { + XELOGE("MetalResolveReadback: failed to create blit encoder"); + do_readback = false; + } else { + blit->copyFromBuffer(source_buffer, source_offset, + rb.buffers[write_index], 0, readback_length); + blit->endEncoding(); + rb.submission_ids[write_index] = GetCurrentSubmission(); + readback_scheduled = true; + } + } + } + + ProcessCompletedSubmissions(); + if (readback_scheduled && use_delayed_sync) { + if (rb.buffers[read_index] == nullptr || + readback_length > rb.sizes[read_index] || + rb.submission_ids[read_index] == 0 || + rb.submission_ids[read_index] > submission_completed_processed_) { + is_cache_miss = true; + read_index = write_index; + } + } + + wait_for_completion = !use_delayed_sync || is_cache_miss; + should_copy = + (readback_mode == ReadbackResolveMode::kSome) ? is_cache_miss : true; + if (readback_scaled && tile_count == 0) { + should_copy = false; + wait_for_completion = false; + } + } + } // Submit the command buffer without waiting - the resolve writes are now // ordered in the same submission as the preceding draws. - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; + uint64_t submission_to_wait = 0; + bool defer_submission = + readback_scheduled && use_delayed_sync && !wait_for_completion; + if (readback_scheduled && wait_for_completion) { + submission_to_wait = GetCurrentSubmission(); + } + if (!defer_submission) { + EndSubmission(false); + if (submission_to_wait) { + CheckSubmissionCompletion(submission_to_wait); + } + } + + if (readback_scheduled && should_copy && readback_buffer && + readback_buffer->buffers[read_index]) { + static uint32_t readback_log_count = 0; + if (readback_log_count < 8) { + ++readback_log_count; + XELOGI( + "MetalResolveReadback: addr=0x{:08X} len={} scaled={} gpu={} " + "scale={}x{} pix_log2={} tiles={} src_off={} " + "scaled_off={} src_len={}", + written_address, written_length, readback_scaled ? 1 : 0, + readback_scaled_gpu ? 1 : 0, scale_x, scale_y, pixel_size_log2, + tile_count, source_offset_bytes_log, scaled_range_offset_bytes, + source_length); + } + const uint8_t* readback_bytes = static_cast( + readback_buffer->buffers[read_index]->contents()); + uint8_t* dest_ptr = memory_->TranslatePhysical(written_address); + if (readback_bytes && dest_ptr) { + if (readback_scaled) { + if (readback_scaled_gpu) { + memory::vastcpy(dest_ptr, const_cast(readback_bytes), + written_length); + } else { + const uint8_t* readback_base = readback_bytes; + if (readback_base_offset_bytes < readback_length) { + readback_base += readback_base_offset_bytes; + } + DownscaleResolveTileData(readback_base, dest_ptr, tile_count, + pixel_size_log2, scale_x, scale_y, + half_pixel_offset); + } + // Scaled resolve data isn't in shared memory; invalidate so CPU memory + // becomes authoritative and shared memory uploads on demand. + if (shared_memory_) { + shared_memory_->MemoryInvalidationCallback(written_address, + written_length, true); + } + if (primitive_processor_) { + primitive_processor_->MemoryInvalidationCallback( + written_address, written_length, true); + } + } else { + memory::vastcpy(dest_ptr, const_cast(readback_bytes), + written_length); + } + } + } + if (readback_scheduled && readback_buffer) { + readback_buffer->current_index = 1u - readback_buffer->current_index; + } return true; } @@ -2990,10 +3616,13 @@ void MetalCommandProcessor::WriteRegister(uint32_t index, uint32_t value) { } MTL::CommandBuffer* MetalCommandProcessor::EnsureCommandBuffer() { - ProcessCompletedSubmissions(); if (current_command_buffer_) { return current_command_buffer_; } + if (!submission_open_) { + XELOGE("EnsureCommandBuffer: no open submission"); + return nullptr; + } if (!command_queue_) { XELOGE("EnsureCommandBuffer: no command queue"); return nullptr; @@ -3007,34 +3636,29 @@ MTL::CommandBuffer* MetalCommandProcessor::EnsureCommandBuffer() { XELOGE("EnsureCommandBuffer: failed to create command buffer"); return nullptr; } - ++submission_current_; current_command_buffer_->retain(); current_command_buffer_->setLabel( NS::String::string("XeniaCommandBuffer", NS::UTF8StringEncoding)); - current_command_buffer_->addCompletedHandler([this](MTL::CommandBuffer*) { - completed_command_buffers_.fetch_add(1, std::memory_order_relaxed); - }); - - if (primitive_processor_) { - primitive_processor_->BeginSubmission(); - } - if (texture_cache_) { - texture_cache_->BeginSubmission(submission_current_); - } - if (primitive_processor_ && !frame_open_) { - primitive_processor_->BeginFrame(); - if (texture_cache_) { - texture_cache_->BeginFrame(); - } - frame_open_ = true; - } return current_command_buffer_; } void MetalCommandProcessor::ProcessCompletedSubmissions() { + CheckSubmissionCompletion(0); +} + +void MetalCommandProcessor::CheckSubmissionCompletion( + uint64_t await_submission) { + if (!completion_timeline_) { + return; + } + if (await_submission) { + completion_timeline_->AwaitSubmissionAndUpdateCompleted(await_submission); + } else { + completion_timeline_->UpdateCompletedSubmission(); + } const uint64_t completed = - completed_command_buffers_.load(std::memory_order_relaxed); + completion_timeline_->GetCompletedSubmissionFromLastUpdate(); if (completed <= submission_completed_processed_) { return; } @@ -3047,6 +3671,108 @@ void MetalCommandProcessor::ProcessCompletedSubmissions() { } } +bool MetalCommandProcessor::BeginSubmission(bool is_guest_command) { + bool is_opening_frame = is_guest_command && !frame_open_; + if (submission_open_ && !is_opening_frame) { + return true; + } + + MaybeStartCapture(); + + uint64_t await_submission = 0; + if (is_opening_frame) { + await_submission = closed_frame_submissions_[frame_current_ % kQueueFrames]; + } + CheckSubmissionCompletion(await_submission); + + if (is_opening_frame) { + frame_completed_ = + std::max(frame_current_, uint64_t(kQueueFrames)) - kQueueFrames; + for (uint64_t frame = frame_completed_ + 1; frame < frame_current_; + ++frame) { + if (closed_frame_submissions_[frame % kQueueFrames] > + GetCompletedSubmission()) { + break; + } + frame_completed_ = frame; + } + } + + if (!submission_open_) { + submission_open_ = true; + if (!EnsureCommandBuffer()) { + submission_open_ = false; + return false; + } + if (primitive_processor_) { + primitive_processor_->BeginSubmission(); + } + if (texture_cache_) { + texture_cache_->BeginSubmission(GetCurrentSubmission()); + } + } + + if (is_opening_frame) { + frame_open_ = true; + if (primitive_processor_) { + primitive_processor_->BeginFrame(); + } + if (texture_cache_) { + texture_cache_->BeginFrame(); + } + if (render_target_cache_) { + render_target_cache_->BeginFrame(); + } + } + + return true; +} + +bool MetalCommandProcessor::EndSubmission(bool is_swap) { + bool is_closing_frame = is_swap && frame_open_; + + if (is_closing_frame) { + if (primitive_processor_) { + primitive_processor_->EndFrame(); + } + } + + if (submission_open_) { + EndRenderEncoder(); + + if (!current_command_buffer_) { + XELOGW("MetalCommandProcessor::EndSubmission: missing command buffer"); + } else { + if (completion_timeline_) { + completion_timeline_->SignalAndAdvance(current_command_buffer_); + } + ScheduleDrawRingRelease(current_command_buffer_); + current_command_buffer_->commit(); + current_command_buffer_->release(); + current_command_buffer_ = nullptr; + } + SetActiveDrawRing(nullptr); + current_draw_index_ = 0; + + submission_open_ = false; + DrainCommandBufferAutoreleasePool(); + } + + if (is_closing_frame) { + if (shared_memory_ && ::cvars::clear_memory_page_state) { + shared_memory_->SetSystemPageBlocksValidWithGpuDataWritten(); + } + frame_open_ = false; + closed_frame_submissions_[frame_current_++ % kQueueFrames] = + GetCurrentSubmission() - 1; + EvictOldReadbackBuffers(readback_buffers_); + } + + return true; +} + +bool MetalCommandProcessor::CanEndSubmissionImmediately() const { return true; } + void MetalCommandProcessor::EnsureCommandBufferAutoreleasePool() { if (command_buffer_autorelease_pool_) { return; @@ -3176,11 +3902,15 @@ void MetalCommandProcessor::BeginCommandBuffer() { } } + bool render_pass_dirty = render_target_cache_ && + render_target_cache_->IsRenderPassDescriptorDirty(); + // If the render pass configuration has changed since the current render // encoder was created (e.g. dummy RT0 -> real RTs, depth/stencil binding), // restart the render encoder with the updated descriptor. if (current_render_encoder_ && - current_render_pass_descriptor_ != pass_descriptor) { + (current_render_pass_descriptor_ != pass_descriptor || + render_pass_dirty)) { EndRenderEncoder(); } @@ -3201,24 +3931,39 @@ void MetalCommandProcessor::BeginCommandBuffer() { UseRenderEncoderAttachmentHeaps(pass_descriptor); } - // Derive viewport/scissor from the actual bound render target rather than - // a hard-coded 1280x720. Prefer color RT 0 from the MetalRenderTargetCache, - // falling back to depth (depth-only passes) and then legacy - // render_target_width_/height_ when needed. + // Derive viewport/scissor from the actual bound render pass attachments. uint32_t rt_width = render_target_width_; uint32_t rt_height = render_target_height_; - if (render_target_cache_) { - MTL::Texture* pass_size_texture = render_target_cache_->GetColorTarget(0); + MTL::Texture* pass_size_texture = nullptr; + if (pass_descriptor) { + if (auto* color_attachments = pass_descriptor->colorAttachments()) { + if (auto* attachment = color_attachments->object(0)) { + pass_size_texture = attachment->texture(); + } + } + if (!pass_size_texture) { + if (auto* depth_attachment = pass_descriptor->depthAttachment()) { + pass_size_texture = depth_attachment->texture(); + } + } + if (!pass_size_texture) { + if (auto* stencil_attachment = pass_descriptor->stencilAttachment()) { + pass_size_texture = stencil_attachment->texture(); + } + } + } + if (!pass_size_texture && render_target_cache_) { + pass_size_texture = render_target_cache_->GetColorTarget(0); if (!pass_size_texture) { pass_size_texture = render_target_cache_->GetDepthTarget(); } if (!pass_size_texture) { pass_size_texture = render_target_cache_->GetDummyColorTarget(); } - if (pass_size_texture) { - rt_width = static_cast(pass_size_texture->width()); - rt_height = static_cast(pass_size_texture->height()); - } + } + if (pass_size_texture) { + rt_width = static_cast(pass_size_texture->width()); + rt_height = static_cast(pass_size_texture->height()); } // Set viewport @@ -3248,19 +3993,7 @@ void MetalCommandProcessor::EnsureDrawRingCapacity() { current_draw_index_ = 0; } -void MetalCommandProcessor::EndCommandBuffer() { - EndRenderEncoder(); - - if (current_command_buffer_) { - ScheduleDrawRingRelease(current_command_buffer_); - current_command_buffer_->commit(); - current_command_buffer_->release(); - current_command_buffer_ = nullptr; - SetActiveDrawRing(nullptr); - current_draw_index_ = 0; - } - DrainCommandBufferAutoreleasePool(); -} +void MetalCommandProcessor::EndCommandBuffer() { EndSubmission(false); } void MetalCommandProcessor::ApplyDepthStencilState( bool primitive_polygonal, reg::RB_DEPTHCONTROL normalized_depth_control) { diff --git a/src/xenia/gpu/metal/metal_command_processor.h b/src/xenia/gpu/metal/metal_command_processor.h index 067ed80cf..a41f1a4c8 100644 --- a/src/xenia/gpu/metal/metal_command_processor.h +++ b/src/xenia/gpu/metal/metal_command_processor.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -10,6 +10,7 @@ #ifndef XENIA_GPU_METAL_METAL_COMMAND_PROCESSOR_H_ #define XENIA_GPU_METAL_METAL_COMMAND_PROCESSOR_H_ +#include #include #include #include @@ -41,6 +42,14 @@ class Heap; class SharedEvent; } // namespace MTL +namespace xe { +namespace ui { +namespace metal { +class MetalGPUCompletionTimeline; +} // namespace metal +} // namespace ui +} // namespace xe + namespace xe { namespace gpu { namespace metal { @@ -78,7 +87,19 @@ class MetalCommandProcessor : public CommandProcessor { MTL::CommandBuffer* GetCurrentCommandBuffer() const { return current_command_buffer_; } + + // Debug marker methods - public so subsystems can annotate their operations. + void UpdateDebugMarkersEnabled(); + void PushDebugMarker(const char* format, ...); + void PopDebugMarker(); + void InsertDebugMarker(const char* format, ...); + bool debug_markers_enabled() const { return debug_markers_enabled_; } + void RequestCapture(); uint32_t current_draw_index() const { return current_draw_index_; } + uint64_t GetCurrentSubmission() const; + uint64_t GetCompletedSubmission() const; + uint64_t GetCurrentFrame() const { return frame_current_; } + uint64_t GetCompletedFrame() const { return frame_completed_; } MTL::CommandBuffer* EnsureCommandBuffer(); void EndRenderEncoder(); void ResetRenderEncoderResourceUsage(); @@ -116,6 +137,7 @@ class MetalCommandProcessor : public CommandProcessor { void IssueSwap(uint32_t frontbuffer_ptr, uint32_t frontbuffer_width, uint32_t frontbuffer_height) override; + void OnPrimaryBufferEnd() override; Shader* LoadShader(xenos::ShaderType shader_type, uint32_t guest_address, const uint32_t* host_address, @@ -143,6 +165,12 @@ class MetalCommandProcessor : public CommandProcessor { void BeginCommandBuffer(); void EndCommandBuffer(); void ProcessCompletedSubmissions(); + void CheckSubmissionCompletion(uint64_t await_submission); + bool BeginSubmission(bool is_guest_command); + bool EndSubmission(bool is_swap); + void MaybeStartCapture(); + void StopCaptureIfActive(); + bool CanEndSubmissionImmediately() const; void EnsureDrawRingCapacity(); void UseRenderEncoderAttachmentHeaps(MTL::RenderPassDescriptor* descriptor); void UseRenderEncoderHeap(MTL::Heap* heap); @@ -354,7 +382,6 @@ class MetalCommandProcessor : public CommandProcessor { // Shared memory for Xbox 360 memory access std::unique_ptr shared_memory_; std::unique_ptr primitive_processor_; - bool frame_open_ = false; bool saw_swap_ = false; uint32_t last_swap_ptr_ = 0; @@ -476,10 +503,24 @@ class MetalCommandProcessor : public CommandProcessor { MTL::BinaryArchive* pipeline_binary_archive_ = nullptr; bool pipeline_binary_archive_dirty_ = false; - std::atomic completed_command_buffers_{0}; - uint64_t submission_current_ = 0; + static constexpr uint32_t kQueueFrames = 3; + + std::unique_ptr completion_timeline_; + bool submission_open_ = false; uint64_t submission_completed_processed_ = 0; + bool frame_open_ = false; + uint64_t frame_current_ = 1; + uint64_t frame_completed_ = 0; + uint64_t closed_frame_submissions_[kQueueFrames] = {}; + + enum class DebugMarkerTarget { + kCommandBuffer, + kRenderEncoder, + }; + bool debug_markers_enabled_ = false; + std::vector debug_marker_stack_; + // Draw counter for ring-buffer descriptor heap allocation // Each draw uses a different region of the descriptor heap to avoid // overwriting previous draws' descriptors before GPU execution @@ -491,6 +532,23 @@ class MetalCommandProcessor : public CommandProcessor { bool gamma_ramp_256_entry_table_up_to_date_ = false; bool gamma_ramp_pwl_up_to_date_ = false; + // Resolve downscale compute shader for scaled resolution readback. + MTL::ComputePipelineState* resolve_downscale_pipeline_ = nullptr; + MTL::Buffer* resolve_downscale_buffer_ = nullptr; + uint32_t resolve_downscale_buffer_size_ = 0; + + // Per-resolve double-buffered readback for delayed sync. + struct ReadbackBuffer { + MTL::Buffer* buffers[2] = {nullptr, nullptr}; + uint32_t sizes[2] = {0, 0}; + uint64_t submission_ids[2] = {0, 0}; + uint32_t current_index = 0; + uint64_t last_used_frame = 0; + }; + void EvictOldReadbackBuffers( + std::unordered_map& buffer_map); + std::unordered_map readback_buffers_; + // Track memory regions written by IssueCopy (resolve) during trace playback. // This prevents the trace player from overwriting resolved data with stale @@ -500,6 +558,10 @@ class MetalCommandProcessor : public CommandProcessor { uint32_t length; }; std::vector resolved_memory_ranges_; + + std::atomic capture_requested_{false}; + MTL::CaptureManager* capture_manager_ = nullptr; + bool capture_active_ = false; }; } // namespace metal diff --git a/src/xenia/gpu/metal/metal_geometry_shader.cc b/src/xenia/gpu/metal/metal_geometry_shader.cc index 55fb8c2a8..64e3699b2 100644 --- a/src/xenia/gpu/metal/metal_geometry_shader.cc +++ b/src/xenia/gpu/metal/metal_geometry_shader.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_geometry_shader.h b/src/xenia/gpu/metal/metal_geometry_shader.h index bc39dd232..6b34af90b 100644 --- a/src/xenia/gpu/metal/metal_geometry_shader.h +++ b/src/xenia/gpu/metal/metal_geometry_shader.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_graphics_system.cc b/src/xenia/gpu/metal/metal_graphics_system.cc index 04302f008..148d90036 100644 --- a/src/xenia/gpu/metal/metal_graphics_system.cc +++ b/src/xenia/gpu/metal/metal_graphics_system.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_graphics_system.h b/src/xenia/gpu/metal/metal_graphics_system.h index 8a114ce12..d0c394dc2 100644 --- a/src/xenia/gpu/metal/metal_graphics_system.h +++ b/src/xenia/gpu/metal/metal_graphics_system.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_heap_pool.cc b/src/xenia/gpu/metal/metal_heap_pool.cc index 04f64bb85..0adcbb98a 100644 --- a/src/xenia/gpu/metal/metal_heap_pool.cc +++ b/src/xenia/gpu/metal/metal_heap_pool.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_heap_pool.h b/src/xenia/gpu/metal/metal_heap_pool.h index b303a8309..f757bd90c 100644 --- a/src/xenia/gpu/metal/metal_heap_pool.h +++ b/src/xenia/gpu/metal/metal_heap_pool.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_primitive_processor.cc b/src/xenia/gpu/metal/metal_primitive_processor.cc index 0f78ad717..562eb182d 100644 --- a/src/xenia/gpu/metal/metal_primitive_processor.cc +++ b/src/xenia/gpu/metal/metal_primitive_processor.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2024 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_primitive_processor.h b/src/xenia/gpu/metal/metal_primitive_processor.h index 539d45b38..4c5d0899b 100644 --- a/src/xenia/gpu/metal/metal_primitive_processor.h +++ b/src/xenia/gpu/metal/metal_primitive_processor.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2024 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_render_target_cache.cc b/src/xenia/gpu/metal/metal_render_target_cache.cc index 40a5aac1b..d8e1f863a 100644 --- a/src/xenia/gpu/metal/metal_render_target_cache.cc +++ b/src/xenia/gpu/metal/metal_render_target_cache.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -24,6 +24,7 @@ #include "xenia/base/assert.h" #include "xenia/base/byte_order.h" #include "xenia/base/logging.h" +#include "xenia/base/math.h" #include "xenia/gpu/draw_util.h" #include "xenia/gpu/gpu_flags.h" #include "xenia/gpu/metal/metal_heap_pool.h" @@ -59,6 +60,18 @@ DEFINE_bool( metal_allow_gamma_unorm16, false, "Allow gamma_render_target_as_unorm16 on Metal despite known issues", "GPU"); +DEFINE_bool(metal_transfer_fast_divmod, true, + "Use fast exact div/mod in Metal transfer shaders", "GPU"); +DEFINE_bool(metal_transfer_tile_instancing, true, + "Use per-tile instanced draws for Metal transfer shaders", "GPU"); +DEFINE_bool( + metal_transfer_msaa_sample_id, true, + "Use sample_id in Metal transfer shaders for MSAA (sample-rate shading)", + "GPU"); +DEFINE_int32(metal_memory_log_rate, 0, + "Log Metal render target/pipeline/instance buffer sizes every N " + "frames (0 to disable)", + "GPU"); namespace xe { namespace gpu { @@ -508,11 +521,30 @@ struct TransferShaderConstants { uint32_t msaa_2x_supported; uint32_t tile_width_samples; uint32_t tile_height_samples; + uint32_t dest_tile_width_pixels; + uint32_t dest_tile_height_pixels; + float dest_tile_width_pixels_inv; + float dest_tile_height_pixels_inv; + float source_pitch_tiles_inv; + float host_depth_source_pitch_tiles_inv; + float dest_pixel_to_ndc_x; + float dest_pixel_to_ndc_y; uint32_t dest_sample_id; uint32_t stencil_mask; uint32_t stencil_clear; }; +struct TransferTileInstance { + float origin_x; + float origin_y; + uint32_t tile_index; + uint32_t padding; + uint32_t source_base_x; + uint32_t source_base_y; + uint32_t host_base_x; + uint32_t host_base_y; +}; + struct TransferClearColorFloatConstants { float color[4]; }; @@ -602,6 +634,10 @@ uint32_t MetalRenderTargetCache::GetMetalEdramDumpFormat(RenderTargetKey key) { // MetalRenderTarget implementation MetalRenderTargetCache::MetalRenderTarget::~MetalRenderTarget() { + if (stencil_view_) { + stencil_view_->release(); + stencil_view_ = nullptr; + } if (draw_texture_ && draw_texture_ != texture_) { draw_texture_->release(); draw_texture_ = nullptr; @@ -749,6 +785,12 @@ void MetalRenderTargetCache::Shutdown(bool from_destructor) { } } transfer_pipelines_.clear(); + for (auto& it : transfer_tile_pipelines_) { + if (it.second) { + it.second->release(); + } + } + transfer_tile_pipelines_.clear(); for (auto& it : edram_load_pipelines_) { if (it.second) { it.second->release(); @@ -799,6 +841,22 @@ void MetalRenderTargetCache::Shutdown(bool from_destructor) { transfer_dummy_buffer_->release(); transfer_dummy_buffer_ = nullptr; } + for (auto& buffer : transfer_tile_instance_buffers_) { + if (buffer) { + buffer->release(); + buffer = nullptr; + } + } + for (auto& retired_list : transfer_tile_instance_retired_buffers_) { + for (auto* buffer : retired_list) { + if (buffer) { + buffer->release(); + } + } + retired_list.clear(); + } + transfer_tile_instance_buffer_sizes_.fill(0); + transfer_tile_instance_buffer_offset_ = 0; for (size_t i = 0; i < xe::countof(transfer_dummy_color_float_); ++i) { if (transfer_dummy_color_float_[i]) { transfer_dummy_color_float_[i]->release(); @@ -828,6 +886,7 @@ void MetalRenderTargetCache::Shutdown(bool from_destructor) { // Destroy all render targets DestroyAllRenderTargets(!from_destructor); + render_target_map_.clear(); if (render_target_heap_pool_) { render_target_heap_pool_->Shutdown(); @@ -1021,10 +1080,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1141,11 +1224,18 @@ kernel void edram_dump_color_32bpp_1xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -1157,8 +1247,10 @@ kernel void edram_dump_color_32bpp_1xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_coord = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -1216,10 +1308,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1336,11 +1452,18 @@ kernel void edram_dump_color_32bpp_2xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -1352,8 +1475,10 @@ kernel void edram_dump_color_32bpp_2xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -1414,10 +1539,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1534,11 +1683,18 @@ kernel void edram_dump_color_32bpp_4xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -1550,8 +1706,10 @@ kernel void edram_dump_color_32bpp_4xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -1614,10 +1772,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1663,11 +1845,18 @@ kernel void edram_dump_depth_32bpp_4xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint2 edram_sample_in_tile = sample_in_tile; uint tile_width_half = tile_size.x >> 1u; edram_sample_in_tile.x = @@ -1686,8 +1875,10 @@ kernel void edram_dump_depth_32bpp_4xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -1759,10 +1950,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1808,11 +2023,18 @@ kernel void edram_dump_depth_32bpp_2xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint2 edram_sample_in_tile = sample_in_tile; uint tile_width_half = tile_size.x >> 1u; edram_sample_in_tile.x = @@ -1831,8 +2053,10 @@ kernel void edram_dump_depth_32bpp_2xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -1902,10 +2126,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -1951,11 +2199,18 @@ kernel void edram_dump_depth_32bpp_1xmsaa( uint3 tid [[thread_position_in_grid]]) { const uint kEdramTileCount = 2048u; - uint2 tile_size = uint2(80u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint2 edram_sample_in_tile = sample_in_tile; uint tile_width_half = tile_size.x >> 1u; edram_sample_in_tile.x = @@ -1974,8 +2229,10 @@ kernel void edram_dump_depth_32bpp_1xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_coord = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -2043,10 +2300,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -2119,11 +2400,18 @@ kernel void edram_dump_color_64bpp_1xmsaa( const uint kEdramTileCount = 2048u; // 64bpp: 40 samples wide per tile instead of 80. - uint2 tile_size = uint2(40u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -2135,8 +2423,10 @@ kernel void edram_dump_color_64bpp_1xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_coord = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -2187,10 +2477,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -2263,11 +2577,18 @@ kernel void edram_dump_color_64bpp_2xmsaa( const uint kEdramTileCount = 2048u; // 64bpp: 40 samples wide per tile instead of 80. - uint2 tile_size = uint2(40u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -2279,8 +2600,10 @@ kernel void edram_dump_color_64bpp_2xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -2334,10 +2657,34 @@ struct EdramDumpConstants { uint dest_pitch_tiles; uint source_pitch_tiles; uint2 resolution_scale; + uint tile_size_x; + uint tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint format; uint flags; + uint padding; }; +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + if (w == 0u) { + q = 0u; + r = 0u; + return; + } + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + constant uint kDumpFormatColorRGBA8 = 0; constant uint kDumpFormatColorRGB10A2Unorm = 1; constant uint kDumpFormatColorRGB10A2Float = 2; @@ -2410,11 +2757,18 @@ kernel void edram_dump_color_64bpp_4xmsaa( const uint kEdramTileCount = 2048u; // 64bpp: 40 samples wide per tile instead of 80. - uint2 tile_size = uint2(40u * constants.resolution_scale.x, - 16u * constants.resolution_scale.y); + uint2 tile_size = uint2(constants.tile_size_x, constants.tile_size_y); - uint2 tile_coord = tid.xy / tile_size; - uint2 sample_in_tile = tid.xy % tile_size; + uint tile_coord_x = 0u; + uint tile_coord_y = 0u; + uint sample_in_tile_x = 0u; + uint sample_in_tile_y = 0u; + XeFastDivMod(tid.x, tile_size.x, constants.tile_size_inv_x, tile_coord_x, + sample_in_tile_x); + XeFastDivMod(tid.y, tile_size.y, constants.tile_size_inv_y, tile_coord_y, + sample_in_tile_y); + uint2 tile_coord = uint2(tile_coord_x, tile_coord_y); + uint2 sample_in_tile = uint2(sample_in_tile_x, sample_in_tile_y); uint rect_tile_index = tile_coord.y * constants.dest_pitch_tiles + tile_coord.x; @@ -2426,8 +2780,10 @@ kernel void edram_dump_color_64bpp_4xmsaa( uint edram_index = wrapped_tile * tile_samples + sample_index; uint source_linear_tile = nonwrapped_tile - constants.source_base_tiles; - uint source_tile_y = source_linear_tile / constants.source_pitch_tiles; - uint source_tile_x = source_linear_tile % constants.source_pitch_tiles; + uint source_tile_y = 0u; + uint source_tile_x = 0u; + XeFastDivMod(source_linear_tile, constants.source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_y, source_tile_x); uint2 source_sample = uint2(source_tile_x * tile_size.x + sample_in_tile.x, source_tile_y * tile_size.y + sample_in_tile.y); @@ -2615,6 +2971,7 @@ void MetalRenderTargetCache::ClearCache() { cleared_render_targets_this_frame_.clear(); dummy_color_targets_.clear(); dummy_color_target_ = nullptr; + render_target_map_.clear(); // Call base implementation RenderTargetCache::ClearCache(); @@ -2628,6 +2985,19 @@ void MetalRenderTargetCache::BeginFrame() { // Call base implementation RenderTargetCache::BeginFrame(); + + if (::cvars::metal_memory_log_rate > 0 && + (frame_id_ % uint64_t(::cvars::metal_memory_log_rate)) == 0) { + XELOGI( + "Metal mem: frame={} rt={} map={} dummy={} pipelines={} " + "tile_pipelines={} inst_buf_sizes=[{}, {}, {}]", + frame_id_, render_targets().size(), render_target_map_.size(), + dummy_color_targets_.size(), transfer_pipelines_.size(), + transfer_tile_pipelines_.size(), + transfer_tile_instance_buffer_sizes_[0], + transfer_tile_instance_buffer_sizes_[1], + transfer_tile_instance_buffer_sizes_[2]); + } } bool MetalRenderTargetCache::Update( @@ -2641,6 +3011,22 @@ bool MetalRenderTargetCache::Update( return false; } + if (::cvars::metal_memory_log_rate > 0) { + static uint64_t memory_log_counter = 0; + if ((++memory_log_counter % uint64_t(::cvars::metal_memory_log_rate)) == + 0) { + XELOGI( + "Metal mem: frame={} rt={} map={} dummy={} pipelines={} " + "tile_pipelines={} inst_buf_sizes=[{}, {}, {}]", + frame_id_, render_targets().size(), render_target_map_.size(), + dummy_color_targets_.size(), transfer_pipelines_.size(), + transfer_tile_pipelines_.size(), + transfer_tile_instance_buffer_sizes_[0], + transfer_tile_instance_buffer_sizes_[1], + transfer_tile_instance_buffer_sizes_[2]); + } + } + // After base class update, retrieve the actual render targets that were // selected This is the KEY to connecting base class management with // Metal-specific rendering @@ -3069,11 +3455,39 @@ MTL::PixelFormat MetalRenderTargetCache::GetDepthPixelFormat( } } +MTL::Texture* MetalRenderTargetCache::GetStencilTextureView( + MetalRenderTarget* render_target) { + if (!render_target) { + return nullptr; + } + if (render_target->stencil_view()) { + return render_target->stencil_view(); + } + RenderTargetKey key = render_target->key(); + if (!key.is_depth) { + return nullptr; + } + MTL::Texture* depth_texture = render_target->texture(); + if (!depth_texture) { + return nullptr; + } + MTL::Texture* view = + depth_texture->newTextureView(MTL::PixelFormatX32_Stencil8); + if (view) { + render_target->SetStencilView(view); + } + return view; +} + MTL::RenderPassDescriptor* MetalRenderTargetCache::GetRenderPassDescriptor( uint32_t expected_sample_count) { - if (!render_pass_descriptor_dirty_ && cached_render_pass_descriptor_) { + if (!render_pass_descriptor_dirty_ && cached_render_pass_descriptor_ && + cached_render_pass_descriptor_sample_count_ == expected_sample_count) { return cached_render_pass_descriptor_; } + if (cached_render_pass_descriptor_sample_count_ != expected_sample_count) { + render_pass_descriptor_dirty_ = true; + } // Release old descriptor if (cached_render_pass_descriptor_) { @@ -3089,6 +3503,7 @@ MTL::RenderPassDescriptor* MetalRenderTargetCache::GetRenderPassDescriptor( return nullptr; } cached_render_pass_descriptor_->retain(); + cached_render_pass_descriptor_sample_count_ = expected_sample_count; bool has_any_render_target = false; bool has_any_color_target = false; @@ -3253,7 +3668,9 @@ MTL::RenderPassDescriptor* MetalRenderTargetCache::GetRenderPassDescriptor( dummy_rt_key.key = 0; dummy_rt_key.is_depth = 0; dummy_rt_key.resource_format = uint32_t(fmt); - dummy_rt_key.msaa_samples = xenos::MsaaSamples(samples); + dummy_rt_key.msaa_samples = (samples >= 4u) ? xenos::MsaaSamples::k4X + : (samples == 2u) ? xenos::MsaaSamples::k2X + : xenos::MsaaSamples::k1X; entry.target = std::make_unique(dummy_rt_key); entry.last_cleared_frame = frame_id_ - 1; MTL::Texture* tex = CreateColorTexture(width, height, fmt, samples); @@ -3301,15 +3718,8 @@ MTL::RenderPassDescriptor* MetalRenderTargetCache::GetRenderPassDescriptor( auto* color_attachment = cached_render_pass_descriptor_->colorAttachments()->object(0); color_attachment->setTexture(dummy_color_target_->draw_texture()); - if (entry.last_cleared_frame != frame_id_) { - color_attachment->setLoadAction(MTL::LoadActionClear); - color_attachment->setClearColor( - MTL::ClearColor::Make(0.0, 0.0, 0.0, 0.0)); - entry.last_cleared_frame = frame_id_; - } else { - color_attachment->setLoadAction(MTL::LoadActionLoad); - } - color_attachment->setStoreAction(MTL::StoreActionStore); + color_attachment->setLoadAction(MTL::LoadActionDontCare); + color_attachment->setStoreAction(MTL::StoreActionDontCare); has_any_render_target = true; if (!coverage_width && dummy_color_target_->draw_texture()) { @@ -3587,8 +3997,14 @@ void MetalRenderTargetCache::DumpRenderTargets( uint32_t source_pitch_tiles; uint32_t resolution_scale_x; uint32_t resolution_scale_y; + uint32_t tile_size_x; + uint32_t tile_size_y; + float tile_size_inv_x; + float tile_size_inv_y; + float source_pitch_tiles_inv; uint32_t format; uint32_t flags; + uint32_t padding; }; MTL::CommandQueue* queue = command_processor_.GetMetalCommandQueue(); @@ -3653,7 +4069,7 @@ void MetalRenderTargetCache::DumpRenderTargets( ::cvars::depth_float24_round) { dump_flags |= kMetalEdramDumpFlagDepthRound; } - stencil_tex = tex->newTextureView(MTL::PixelFormatX32_Stencil8); + stencil_tex = GetStencilTextureView(rt); if (stencil_tex) { dump_flags |= kMetalEdramDumpFlagHasStencil; } @@ -3734,9 +4150,6 @@ void MetalRenderTargetCache::DumpRenderTargets( "key=0x{:08X} (is_depth={}, is_64bpp={}, msaa={})", key.key, key.is_depth ? 1 : 0, is_64bpp ? 1 : 0, static_cast(key.msaa_samples)); - if (stencil_tex) { - stencil_tex->release(); - } continue; } @@ -3752,9 +4165,6 @@ void MetalRenderTargetCache::DumpRenderTargets( uint32_t dispatch_count = rect.GetDispatches(dump_pitch, dump_row_length_used, dispatches); if (!dispatch_count) { - if (stencil_tex) { - stencil_tex->release(); - } continue; } @@ -3768,8 +4178,21 @@ void MetalRenderTargetCache::DumpRenderTargets( constants.source_pitch_tiles = key.GetPitchTiles(); constants.resolution_scale_x = scale_x; constants.resolution_scale_y = scale_y; + uint32_t tile_size_x = (is_64bpp ? 40u : 80u) * scale_x; + uint32_t tile_size_y = 16u * scale_y; + constants.tile_size_x = tile_size_x; + constants.tile_size_y = tile_size_y; + constants.tile_size_inv_x = + tile_size_x ? (1.0f / float(tile_size_x)) : 0.0f; + constants.tile_size_inv_y = + tile_size_y ? (1.0f / float(tile_size_y)) : 0.0f; + constants.source_pitch_tiles_inv = + constants.source_pitch_tiles + ? (1.0f / float(constants.source_pitch_tiles)) + : 0.0f; constants.format = dump_format; constants.flags = dump_flags; + constants.padding = 0; encoder->setComputePipelineState(dump_pipeline); encoder->setTexture(tex, 0); @@ -3796,10 +4219,6 @@ void MetalRenderTargetCache::DumpRenderTargets( MTL::Size threadgroups = MTL::Size::Make(groups_x, groups_y, 1); encoder->dispatchThreadgroups(threadgroups, threads_per_group); } - - if (stencil_tex) { - stencil_tex->release(); - } } encoder->endEncoding(); @@ -4639,6 +5058,8 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( } uint32_t dest_sample_count = MsaaSamplesToCount(dest_key.msaa_samples); + bool transfer_use_sample_id_default = + dest_sample_count > 1 && ::cvars::metal_transfer_msaa_sample_id; uint32_t dest_width = uint32_t(dest_texture->width()); uint32_t dest_height = uint32_t(dest_texture->height()); @@ -4692,6 +5113,235 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( return true; }; + struct TransferTileBatch { + MTL::Buffer* buffer = nullptr; + size_t buffer_offset = 0; + uint32_t instance_count = 0; + MTL::ScissorRect scissor = {}; + }; + + struct TransferTileBatchBuildInfo { + TransferTileBatch batch; + uint32_t tile_x_start = 0; + uint32_t tile_x_end = 0; + uint32_t tile_y_start = 0; + uint32_t tile_y_end = 0; + }; + + auto allocate_instance_buffer = [&](size_t size, MTL::Buffer*& buffer, + size_t& offset) -> bool { + if (!device_) { + return false; + } + uint32_t buffer_index = + uint32_t(frame_id_ % kTransferInstanceBufferCount); + if (transfer_tile_instance_buffer_frame_id_ != frame_id_) { + transfer_tile_instance_buffer_frame_id_ = frame_id_; + transfer_tile_instance_buffer_offset_ = 0; + auto& retired_buffers = + transfer_tile_instance_retired_buffers_[buffer_index]; + for (auto* retired_buffer : retired_buffers) { + if (retired_buffer) { + retired_buffer->release(); + } + } + retired_buffers.clear(); + } + constexpr size_t kAlignment = 256; + size_t aligned_offset = + xe::align(transfer_tile_instance_buffer_offset_, size_t(kAlignment)); + size_t required = aligned_offset + size; + if (!transfer_tile_instance_buffers_[buffer_index] || + transfer_tile_instance_buffer_sizes_[buffer_index] < required) { + size_t new_size = xe::round_up(required, 65536); + if (transfer_tile_instance_buffers_[buffer_index]) { + transfer_tile_instance_retired_buffers_[buffer_index].push_back( + transfer_tile_instance_buffers_[buffer_index]); + transfer_tile_instance_buffers_[buffer_index] = nullptr; + } + MTL::ResourceOptions options = MTL::ResourceStorageModeShared | + MTL::ResourceCPUCacheModeWriteCombined; + MTL::Buffer* new_buffer = device_->newBuffer(new_size, options); + if (!new_buffer) { + transfer_tile_instance_buffer_sizes_[buffer_index] = 0; + return false; + } + transfer_tile_instance_buffers_[buffer_index] = new_buffer; + transfer_tile_instance_buffer_sizes_[buffer_index] = new_size; + } + buffer = transfer_tile_instance_buffers_[buffer_index]; + if (!buffer) { + return false; + } + offset = aligned_offset; + transfer_tile_instance_buffer_offset_ = aligned_offset + size; + return true; + }; + + auto build_tile_batches = + [&](const Transfer::Rectangle* rectangles, uint32_t rectangle_count, + const TransferShaderConstants& constants, bool uses_host_depth, + bool host_depth_is_copy, + std::vector& out_batches) -> bool { + out_batches.clear(); + if (!constants.dest_tile_width_pixels || + !constants.dest_tile_height_pixels) { + return false; + } + if (!cmd) { + return false; + } + uint32_t max_tile_x = + (dest_width + constants.dest_tile_width_pixels - 1) / + constants.dest_tile_width_pixels; + uint32_t max_tile_y = + (dest_height + constants.dest_tile_height_pixels - 1) / + constants.dest_tile_height_pixels; + if (!max_tile_x || !max_tile_y) { + return false; + } + --max_tile_x; + --max_tile_y; + std::vector build_infos; + size_t total_instance_bytes = 0; + constexpr size_t kAlignment = 256; + for (uint32_t rect_index = 0; rect_index < rectangle_count; + ++rect_index) { + uint32_t scaled_x = 0; + uint32_t scaled_y = 0; + uint32_t scaled_width = 0; + uint32_t scaled_height = 0; + if (!get_scaled_rect(rectangles[rect_index], scaled_x, scaled_y, + scaled_width, scaled_height)) { + continue; + } + uint32_t tile_x_start = scaled_x / constants.dest_tile_width_pixels; + uint32_t tile_y_start = scaled_y / constants.dest_tile_height_pixels; + uint32_t tile_x_end = + (scaled_x + scaled_width - 1) / constants.dest_tile_width_pixels; + uint32_t tile_y_end = + (scaled_y + scaled_height - 1) / constants.dest_tile_height_pixels; + tile_x_end = std::min(tile_x_end, max_tile_x); + tile_y_end = std::min(tile_y_end, max_tile_y); + if (tile_x_start > tile_x_end || tile_y_start > tile_y_end) { + continue; + } + uint32_t tiles_x = tile_x_end - tile_x_start + 1; + uint32_t tiles_y = tile_y_end - tile_y_start + 1; + uint32_t tile_count = tiles_x * tiles_y; + if (!tile_count) { + continue; + } + TransferTileBatchBuildInfo info; + info.tile_x_start = tile_x_start; + info.tile_x_end = tile_x_end; + info.tile_y_start = tile_y_start; + info.tile_y_end = tile_y_end; + info.batch.instance_count = tile_count; + info.batch.scissor.x = scaled_x; + info.batch.scissor.y = scaled_y; + info.batch.scissor.width = scaled_width; + info.batch.scissor.height = scaled_height; + total_instance_bytes = xe::align(total_instance_bytes, kAlignment); + info.batch.buffer_offset = total_instance_bytes; + total_instance_bytes += + size_t(tile_count) * sizeof(TransferTileInstance); + build_infos.push_back(info); + } + + if (build_infos.empty() || !total_instance_bytes) { + return false; + } + + MTL::Buffer* buffer = nullptr; + size_t buffer_base_offset = 0; + if (!allocate_instance_buffer(total_instance_bytes, buffer, + buffer_base_offset)) { + return false; + } + + uint8_t* base_ptr = + reinterpret_cast(buffer->contents()) + buffer_base_offset; + uint32_t source_pitch_tiles = constants.address.source_pitch; + uint32_t source_tile_width_pixels = + tile_width_samples >> + ((constants.source_is_64bpp != 0u) + + (constants.source_msaa_samples >= 4u ? 1u : 0u)); + uint32_t source_tile_height_pixels = + tile_height_samples >> + (constants.source_msaa_samples >= 2u ? 1u : 0u); + uint32_t host_tile_width_pixels = + tile_width_samples >> + (constants.host_depth_source_msaa_samples >= 4u ? 1u : 0u); + uint32_t host_tile_height_pixels = + tile_height_samples >> + (constants.host_depth_source_msaa_samples >= 2u ? 1u : 0u); + + for (const auto& info : build_infos) { + uint8_t* batch_ptr = base_ptr + info.batch.buffer_offset; + auto* instances = reinterpret_cast(batch_ptr); + uint32_t instance_index = 0; + for (uint32_t tile_y = info.tile_y_start; tile_y <= info.tile_y_end; + ++tile_y) { + uint32_t row_base = tile_y * constants.address.dest_pitch; + float origin_y = float(tile_y * constants.dest_tile_height_pixels); + for (uint32_t tile_x = info.tile_x_start; tile_x <= info.tile_x_end; + ++tile_x) { + TransferTileInstance& instance = instances[instance_index++]; + instance.origin_x = + float(tile_x * constants.dest_tile_width_pixels); + instance.origin_y = origin_y; + instance.tile_index = row_base + tile_x; + uint32_t dest_tile_index = instance.tile_index; + uint32_t source_tile_index = + uint32_t(int32_t(dest_tile_index) + + constants.address.source_to_dest) & + (xenos::kEdramTileCount - 1u); + uint32_t source_tile_index_y = 0u; + uint32_t source_tile_index_x = 0u; + if (source_pitch_tiles) { + source_tile_index_y = source_tile_index / source_pitch_tiles; + source_tile_index_x = + source_tile_index - source_tile_index_y * source_pitch_tiles; + } + instance.source_base_x = + source_tile_index_x * source_tile_width_pixels; + instance.source_base_y = + source_tile_index_y * source_tile_height_pixels; + instance.host_base_x = 0; + instance.host_base_y = 0; + if (uses_host_depth && !host_depth_is_copy) { + uint32_t host_pitch_tiles = + constants.host_depth_address.source_pitch; + uint32_t host_tile_index = + uint32_t(int32_t(dest_tile_index) + + constants.host_depth_address.source_to_dest) & + (xenos::kEdramTileCount - 1u); + uint32_t host_tile_index_y = 0u; + uint32_t host_tile_index_x = 0u; + if (host_pitch_tiles) { + host_tile_index_y = host_tile_index / host_pitch_tiles; + host_tile_index_x = + host_tile_index - host_tile_index_y * host_pitch_tiles; + } + instance.host_base_x = host_tile_index_x * host_tile_width_pixels; + instance.host_base_y = + host_tile_index_y * host_tile_height_pixels; + } + } + } + TransferTileBatch batch = info.batch; + batch.buffer = buffer; + batch.buffer_offset = buffer_base_offset + info.batch.buffer_offset; + out_batches.push_back(batch); + } + + if (out_batches.empty()) { + return false; + } + return true; + }; + std::vector filtered_transfers; bool used_blit = false; MTL::BlitCommandEncoder* blit_encoder = nullptr; @@ -4702,74 +5352,221 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( return blit_encoder; }; - // Fast path: when source/dest share identical EDRAM layout and format, + // Fast path: when source/dest share compatible EDRAM layout and format, // use a blit instead of shader-based transfers. - if (!transfers.empty() && !dest_is_depth) { - for (const Transfer& transfer : transfers) { + if (!transfers.empty()) { + auto try_blit_transfer = [&](const Transfer& transfer) -> bool { auto* source_rt = static_cast(transfer.source); if (!source_rt || transfer.host_depth_source) { - filtered_transfers.push_back(transfer); - continue; + return false; } + RenderTargetKey source_key = source_rt->key(); - if (source_key.is_depth) { - filtered_transfers.push_back(transfer); - continue; + if (dest_is_depth != source_key.is_depth) { + return false; } if (source_key.resource_format != dest_key.resource_format || source_key.msaa_samples != dest_key.msaa_samples || - source_key.base_tiles != dest_key.base_tiles || source_key.pitch_tiles_at_32bpp != dest_key.pitch_tiles_at_32bpp) { - filtered_transfers.push_back(transfer); - continue; + return false; } - MTL::Texture* source_texture = source_rt->transfer_texture(); + + bool base_tiles_match = source_key.base_tiles == dest_key.base_tiles; + if (dest_is_depth && !base_tiles_match) { + return false; + } + + MTL::Texture* source_texture = dest_is_depth + ? source_rt->texture() + : source_rt->transfer_texture(); if (!source_texture) { - filtered_transfers.push_back(transfer); - continue; + return false; + } + if (!dest_is_depth) { + MTL::PixelFormat expected_format = + GetColorOwnershipTransferPixelFormat(source_key.GetColorFormat(), + nullptr); + assert_true(source_texture->pixelFormat() == expected_format, + "Transfer source must use ownership pixel format"); } - MTL::PixelFormat expected_format = GetColorOwnershipTransferPixelFormat( - source_key.GetColorFormat(), nullptr); - assert_true(source_texture->pixelFormat() == expected_format, - "Transfer source must use ownership pixel format"); if (source_texture->pixelFormat() != dest_texture->pixelFormat() || + source_texture->sampleCount() != dest_texture->sampleCount() || source_texture->sampleCount() != 1 || - dest_texture->sampleCount() != 1 || source_texture->width() != dest_width || source_texture->height() != dest_height) { - filtered_transfers.push_back(transfer); - continue; + return false; } + Transfer::Rectangle rectangles[Transfer::kMaxRectanglesWithCutout]; uint32_t rectangle_count = transfer.GetRectangles( dest_key.base_tiles, dest_key.pitch_tiles_at_32bpp, dest_key.msaa_samples, IsKey64bpp(dest_key), rectangles, resolve_clear_rectangle); if (!rectangle_count) { - continue; + return false; } + MTL::BlitCommandEncoder* blit = ensure_blit_encoder(); if (!blit) { - filtered_transfers.push_back(transfer); - continue; + return false; } - for (uint32_t rect_index = 0; rect_index < rectangle_count; - ++rect_index) { - uint32_t scaled_x = 0; - uint32_t scaled_y = 0; - uint32_t scaled_width = 0; - uint32_t scaled_height = 0; - if (!get_scaled_rect(rectangles[rect_index], scaled_x, scaled_y, - scaled_width, scaled_height)) { - continue; + + if (base_tiles_match || dest_is_depth) { + for (uint32_t rect_index = 0; rect_index < rectangle_count; + ++rect_index) { + uint32_t scaled_x = 0; + uint32_t scaled_y = 0; + uint32_t scaled_width = 0; + uint32_t scaled_height = 0; + if (!get_scaled_rect(rectangles[rect_index], scaled_x, scaled_y, + scaled_width, scaled_height)) { + continue; + } + MTL::Origin origin = MTL::Origin::Make(scaled_x, scaled_y, 0); + MTL::Size size = MTL::Size::Make(scaled_width, scaled_height, 1); + blit->copyFromTexture(source_texture, 0, 0, origin, size, + dest_texture, 0, 0, origin); + } + } else { + // Base-tile offset blit (color only, non-MSAA, tile-aligned). + uint32_t pitch_tiles = dest_key.pitch_tiles_at_32bpp; + if (!pitch_tiles) { + return false; + } + uint32_t tile_width_pixels = + tile_width_samples >> + ((IsKey64bpp(dest_key) ? 1u : 0u) + + uint32_t(dest_key.msaa_samples >= xenos::MsaaSamples::k4X)); + uint32_t tile_height_pixels = + tile_height_samples >> + uint32_t(dest_key.msaa_samples >= xenos::MsaaSamples::k2X); + if (!tile_width_pixels || !tile_height_pixels) { + return false; + } + uint32_t delta_tiles = (dest_key.base_tiles - source_key.base_tiles) & + (xenos::kEdramTileCount - 1u); + uint32_t delta_rows = delta_tiles / pitch_tiles; + uint32_t delta_x = delta_tiles % pitch_tiles; + uint32_t total_rows = + (xenos::kEdramTileCount + pitch_tiles - 1u) / pitch_tiles; + + struct ScaledRect { + uint32_t x; + uint32_t y; + uint32_t width; + uint32_t height; + }; + std::vector scaled_rects; + scaled_rects.reserve(rectangle_count); + for (uint32_t rect_index = 0; rect_index < rectangle_count; + ++rect_index) { + uint32_t scaled_x = 0; + uint32_t scaled_y = 0; + uint32_t scaled_width = 0; + uint32_t scaled_height = 0; + if (!get_scaled_rect(rectangles[rect_index], scaled_x, scaled_y, + scaled_width, scaled_height)) { + continue; + } + if ((scaled_x % tile_width_pixels) || + (scaled_y % tile_height_pixels) || + (scaled_width % tile_width_pixels) || + (scaled_height % tile_height_pixels)) { + return false; + } + if (!scaled_width || !scaled_height) { + continue; + } + scaled_rects.push_back( + {scaled_x, scaled_y, scaled_width, scaled_height}); + } + if (scaled_rects.empty()) { + return false; + } + + for (const auto& rect : scaled_rects) { + uint32_t tile_x = rect.x / tile_width_pixels; + uint32_t tile_y = rect.y / tile_height_pixels; + uint32_t tiles_w = rect.width / tile_width_pixels; + uint32_t tiles_h = rect.height / tile_height_pixels; + if (!tiles_w || !tiles_h) { + continue; + } + + uint32_t source_tile_x_base = tile_x + delta_x; + uint32_t source_tile_x = source_tile_x_base % pitch_tiles; + uint32_t source_tile_y = + tile_y + delta_rows + (source_tile_x_base / pitch_tiles); + if (source_tile_y >= total_rows) { + source_tile_y %= total_rows; + } + + uint32_t rows_before_wrap = + std::min(tiles_h, total_rows - source_tile_y); + uint32_t rows_after_wrap = tiles_h - rows_before_wrap; + + uint32_t tiles_before_wrap_x = + (source_tile_x + tiles_w <= pitch_tiles) + ? tiles_w + : (pitch_tiles - source_tile_x); + uint32_t tiles_after_wrap_x = tiles_w - tiles_before_wrap_x; + + for (uint32_t wrap_y = 0; wrap_y <= (rows_after_wrap ? 1u : 0u); + ++wrap_y) { + uint32_t y_offset_tiles = wrap_y ? rows_before_wrap : 0u; + uint32_t rows = wrap_y ? rows_after_wrap : rows_before_wrap; + if (!rows) { + continue; + } + uint32_t dest_y_pixels = + rect.y + y_offset_tiles * tile_height_pixels; + uint32_t source_y_tiles = wrap_y ? 0u : source_tile_y; + uint32_t source_y_pixels = source_y_tiles * tile_height_pixels; + uint32_t height_pixels = rows * tile_height_pixels; + + // X segment 0. + if (tiles_before_wrap_x) { + uint32_t dest_x_pixels = rect.x; + uint32_t source_x_pixels = source_tile_x * tile_width_pixels; + uint32_t width_pixels = tiles_before_wrap_x * tile_width_pixels; + MTL::Origin src_origin = + MTL::Origin::Make(source_x_pixels, source_y_pixels, 0); + MTL::Origin dst_origin = + MTL::Origin::Make(dest_x_pixels, dest_y_pixels, 0); + MTL::Size size = + MTL::Size::Make(width_pixels, height_pixels, 1); + blit->copyFromTexture(source_texture, 0, 0, src_origin, size, + dest_texture, 0, 0, dst_origin); + } + + // X segment 1 (wrap). + if (tiles_after_wrap_x) { + uint32_t dest_x_pixels = + rect.x + tiles_before_wrap_x * tile_width_pixels; + uint32_t source_x_pixels = 0; + uint32_t width_pixels = tiles_after_wrap_x * tile_width_pixels; + MTL::Origin src_origin = + MTL::Origin::Make(source_x_pixels, source_y_pixels, 0); + MTL::Origin dst_origin = + MTL::Origin::Make(dest_x_pixels, dest_y_pixels, 0); + MTL::Size size = + MTL::Size::Make(width_pixels, height_pixels, 1); + blit->copyFromTexture(source_texture, 0, 0, src_origin, size, + dest_texture, 0, 0, dst_origin); + } + } } - MTL::Origin origin = MTL::Origin::Make(scaled_x, scaled_y, 0); - MTL::Size size = MTL::Size::Make(scaled_width, scaled_height, 1); - blit->copyFromTexture(source_texture, 0, 0, origin, size, - dest_texture, 0, 0, origin); } + used_blit = true; any_transfers_done = true; + return true; + }; + + for (const Transfer& transfer : transfers) { + if (!try_blit_transfer(transfer)) { + filtered_transfers.push_back(transfer); + } } } @@ -4903,6 +5700,24 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( } } + const TransferModeInfo& mode_info = + kTransferModeInfos[size_t(shader_key.mode)]; + bool transfer_use_sample_id = transfer_use_sample_id_default; + if (transfer_use_sample_id) { + bool source_is_multisample = + source_key.msaa_samples != xenos::MsaaSamples::k1X; + bool host_depth_is_multisample = + mode_info.uses_host_depth && + shader_key.host_depth_source_msaa_samples != + xenos::MsaaSamples::k1X && + !shader_key.host_depth_source_is_copy; + if (!source_is_multisample && !host_depth_is_multisample) { + transfer_use_sample_id = false; + } + } + shader_key.dest_sample_id_from_sample = + transfer_use_sample_id ? 1u : 0u; + transfer_invocations_.emplace_back(transfer, shader_key); if (pass) { transfer_invocations_.back().transfer.host_depth_source = nullptr; @@ -4948,12 +5763,31 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( } MTL::RenderCommandEncoder* encoder = ensure_transfer_encoder(); if (encoder) { + bool transfer_viewport_full_set = false; + auto set_full_transfer_viewport = [&]() { + if (transfer_viewport_full_set) { + return; + } + MTL::Viewport vp; + vp.originX = 0.0; + vp.originY = 0.0; + vp.width = double(dest_width); + vp.height = double(dest_height); + vp.znear = 0.0; + vp.zfar = 1.0; + encoder->setViewport(vp); + transfer_viewport_full_set = true; + }; for (const auto& invocation : transfer_invocations_) { const Transfer& transfer = invocation.transfer; const TransferShaderKey& shader_key = invocation.shader_key; const TransferModeInfo& mode_info = kTransferModeInfos[size_t(shader_key.mode)]; bool is_stencil_bit = mode_info.output == TransferOutput::kStencilBit; + bool needs_source_stencil = + !mode_info.source_is_color && + (mode_info.output == TransferOutput::kColor || + mode_info.output == TransferOutput::kStencilBit); auto* source_rt = static_cast(transfer.source); if (!source_rt) { @@ -4968,14 +5802,6 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( source_key.GetColorFormat(), &source_is_uint); } - MTL::RenderPipelineState* pipeline = GetOrCreateTransferPipelines( - shader_key, dest_pixel_format, dest_is_uint); - if (!pipeline) { - continue; - } - - encoder->setRenderPipelineState(pipeline); - if (is_stencil_bit) { // Depth/stencil state set per-bit below. } else if (dest_is_depth) { @@ -5004,13 +5830,13 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( continue; } encoder->setFragmentTexture(depth_texture, 0); - MTL::Texture* stencil_texture = - depth_texture->newTextureView(MTL::PixelFormatX32_Stencil8); - if (!stencil_texture) { - continue; + if (needs_source_stencil) { + MTL::Texture* stencil_texture = GetStencilTextureView(source_rt); + if (!stencil_texture) { + continue; + } + encoder->setFragmentTexture(stencil_texture, 1); } - encoder->setFragmentTexture(stencil_texture, 1); - stencil_texture->release(); } // Bind host depth source if needed. @@ -5081,6 +5907,39 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( constants.msaa_2x_supported = msaa_2x_supported_ ? 1 : 0; constants.tile_width_samples = tile_width_samples; constants.tile_height_samples = tile_height_samples; + uint32_t dest_tile_width_pixels = + tile_width_samples >> + ((constants.dest_is_64bpp != 0u) + + (constants.dest_msaa_samples >= 4u ? 1u : 0u)); + uint32_t dest_tile_height_pixels = + tile_height_samples >> + (constants.dest_msaa_samples >= 2u ? 1u : 0u); + constants.dest_tile_width_pixels = dest_tile_width_pixels; + constants.dest_tile_height_pixels = dest_tile_height_pixels; + constants.dest_tile_width_pixels_inv = + dest_tile_width_pixels ? (1.0f / float(dest_tile_width_pixels)) + : 0.0f; + constants.dest_tile_height_pixels_inv = + dest_tile_height_pixels ? (1.0f / float(dest_tile_height_pixels)) + : 0.0f; + uint32_t source_pitch_tiles = source_key.GetPitchTiles(); + constants.source_pitch_tiles_inv = + source_pitch_tiles ? (1.0f / float(source_pitch_tiles)) : 0.0f; + constants.host_depth_source_pitch_tiles_inv = 0.0f; + if (mode_info.uses_host_depth && + !shader_key.host_depth_source_is_copy) { + auto* host_depth_rt = + static_cast(transfer.host_depth_source); + if (host_depth_rt) { + uint32_t host_pitch_tiles = host_depth_rt->key().GetPitchTiles(); + constants.host_depth_source_pitch_tiles_inv = + host_pitch_tiles ? (1.0f / float(host_pitch_tiles)) : 0.0f; + } + } + constants.dest_pixel_to_ndc_x = + dest_width ? (2.0f / float(dest_width)) : 0.0f; + constants.dest_pixel_to_ndc_y = + dest_height ? (2.0f / float(dest_height)) : 0.0f; constants.dest_sample_id = 0; Transfer::Rectangle rectangles[Transfer::kMaxRectanglesWithCutout]; @@ -5092,6 +5951,62 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( continue; } + std::vector tile_batches; + bool use_tile_instancing = false; + if (::cvars::metal_transfer_tile_instancing) { + use_tile_instancing = build_tile_batches( + rectangles, rectangle_count, constants, + mode_info.uses_host_depth, + shader_key.host_depth_source_is_copy != 0, tile_batches); + } + + MTL::RenderPipelineState* pipeline = GetOrCreateTransferPipelines( + shader_key, dest_pixel_format, dest_is_uint, use_tile_instancing); + if (!pipeline) { + continue; + } + + encoder->setRenderPipelineState(pipeline); + + bool use_sample_id_for_invocation = + shader_key.dest_sample_id_from_sample != 0; + auto draw_transfer_samples = [&](auto&& draw_fn) { + if (use_sample_id_for_invocation || dest_sample_count <= 1) { + draw_fn(0); + return; + } + for (uint32_t sample_id = 0; sample_id < dest_sample_count; + ++sample_id) { + draw_fn(sample_id); + } + }; + + auto draw_transfer = [&](uint32_t sample_id) { + constants.dest_sample_id = sample_id; + if (use_tile_instancing) { + set_full_transfer_viewport(); + encoder->setVertexBytes(&constants, sizeof(constants), 0); + encoder->setFragmentBytes(&constants, sizeof(constants), 0); + for (const auto& batch : tile_batches) { + encoder->setScissorRect(batch.scissor); + encoder->setVertexBuffer(batch.buffer, batch.buffer_offset, 1); + encoder->drawPrimitives(MTL::PrimitiveTypeTriangleStrip, + NS::UInteger(0), NS::UInteger(4), + NS::UInteger(batch.instance_count)); + } + } else { + encoder->setFragmentBytes(&constants, sizeof(constants), 0); + for (uint32_t rect_index = 0; rect_index < rectangle_count; + ++rect_index) { + if (!set_rect_viewport(encoder, rectangles[rect_index])) { + continue; + } + encoder->drawPrimitives(MTL::PrimitiveTypeTriangle, + NS::UInteger(0), NS::UInteger(3)); + } + } + }; + if (is_stencil_bit) { for (uint32_t bit = 0; bit < 8; ++bit) { MTL::DepthStencilState* stencil_state = @@ -5103,36 +6018,12 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( constants.stencil_clear = 0; encoder->setDepthStencilState(stencil_state); encoder->setStencilReferenceValue(uint32_t(1) << bit); - for (uint32_t sample_id = 0; sample_id < dest_sample_count; - ++sample_id) { - constants.dest_sample_id = sample_id; - encoder->setFragmentBytes(&constants, sizeof(constants), 0); - for (uint32_t rect_index = 0; rect_index < rectangle_count; - ++rect_index) { - if (!set_rect_viewport(encoder, rectangles[rect_index])) { - continue; - } - encoder->drawPrimitives(MTL::PrimitiveTypeTriangle, - NS::UInteger(0), NS::UInteger(3)); - } - } + draw_transfer_samples(draw_transfer); } } else { constants.stencil_mask = 0; constants.stencil_clear = 0; - for (uint32_t sample_id = 0; sample_id < dest_sample_count; - ++sample_id) { - constants.dest_sample_id = sample_id; - encoder->setFragmentBytes(&constants, sizeof(constants), 0); - for (uint32_t rect_index = 0; rect_index < rectangle_count; - ++rect_index) { - if (!set_rect_viewport(encoder, rectangles[rect_index])) { - continue; - } - encoder->drawPrimitives(MTL::PrimitiveTypeTriangle, - NS::UInteger(0), NS::UInteger(3)); - } - } + draw_transfer_samples(draw_transfer); } any_transfers_done = true; } @@ -5337,9 +6228,11 @@ void MetalRenderTargetCache::PerformTransfersAndResolveClears( MTL::RenderPipelineState* MetalRenderTargetCache::GetOrCreateTransferPipelines( const TransferShaderKey& key, MTL::PixelFormat dest_format, - bool dest_is_uint) { - auto it = transfer_pipelines_.find(key); - if (it != transfer_pipelines_.end()) { + bool dest_is_uint, bool tile_instanced) { + auto& pipeline_map = + tile_instanced ? transfer_tile_pipelines_ : transfer_pipelines_; + auto it = pipeline_map.find(key); + if (it != pipeline_map.end()) { return it->second; } @@ -5361,6 +6254,22 @@ MTL::RenderPipelineState* MetalRenderTargetCache::GetOrCreateTransferPipelines( if (source_is_color) { GetColorOwnershipTransferPixelFormat(source_color_format, &source_is_uint); } + bool source_is_64bpp = false; + if (source_is_color) { + source_is_64bpp = + xenos::IsColorRenderTargetFormat64bpp(source_color_format); + } + bool dest_is_depth = output != TransferOutput::kColor; + bool dest_is_64bpp = false; + if (!dest_is_depth) { + dest_is_64bpp = + xenos::IsColorRenderTargetFormat64bpp(dest_color_format) || + (dest_color_format == xenos::ColorRenderTargetFormat::k_8_8_8_8_GAMMA && + gamma_render_target_as_unorm16_); + } + bool source_needs_stencil = + !source_is_color && (output == TransferOutput::kColor || + output == TransferOutput::kStencilBit); uint32_t dest_component_count = 1; if (output == TransferOutput::kColor) { @@ -5388,18 +6297,35 @@ MTL::RenderPipelineState* MetalRenderTargetCache::GetOrCreateTransferPipelines( std::string source; source.reserve(16384); append_define(source, "XE_TRANSFER_SOURCE_IS_COLOR", source_is_color ? 1 : 0); + append_define(source, "XE_TRANSFER_SOURCE_IS_DEPTH", source_is_color ? 0 : 1); + append_define(source, "XE_TRANSFER_SOURCE_NEEDS_STENCIL", + source_needs_stencil ? 1 : 0); append_define(source, "XE_TRANSFER_SOURCE_IS_UINT", source_is_uint ? 1 : 0); + append_define(source, "XE_TRANSFER_SOURCE_IS_64BPP", source_is_64bpp ? 1 : 0); append_define(source, "XE_TRANSFER_SOURCE_IS_MULTISAMPLE", source_is_multisample ? 1 : 0); append_define(source, "XE_TRANSFER_DEST_IS_UINT", dest_is_uint ? 1 : 0); + append_define(source, "XE_TRANSFER_DEST_IS_DEPTH", dest_is_depth ? 1 : 0); + append_define(source, "XE_TRANSFER_DEST_IS_64BPP", dest_is_64bpp ? 1 : 0); append_define(source, "XE_TRANSFER_DEST_COMPONENTS", dest_component_count); append_define(source, "XE_TRANSFER_DEST_IS_MULTISAMPLE", dest_is_multisample ? 1 : 0); + append_define(source, "XE_TRANSFER_DEST_SAMPLE_ID_FROM_SAMPLE", + key.dest_sample_id_from_sample ? 1 : 0); append_define(source, "XE_TRANSFER_HAS_HOST_DEPTH", has_host_depth ? 1 : 0); append_define(source, "XE_TRANSFER_HOST_DEPTH_IS_COPY", host_depth_is_copy ? 1 : 0); append_define(source, "XE_TRANSFER_HOST_DEPTH_IS_MULTISAMPLE", host_depth_is_multisample ? 1 : 0); + append_define(source, "XE_TRANSFER_SOURCE_FORMAT", + key.source_resource_format); + append_define(source, "XE_TRANSFER_DEST_FORMAT", key.dest_resource_format); + append_define(source, "XE_TRANSFER_SOURCE_MSAA_SAMPLES", + MsaaSamplesToCount(key.source_msaa_samples)); + append_define(source, "XE_TRANSFER_DEST_MSAA_SAMPLES", + MsaaSamplesToCount(key.dest_msaa_samples)); + append_define(source, "XE_TRANSFER_HOST_DEPTH_MSAA_SAMPLES", + MsaaSamplesToCount(key.host_depth_source_msaa_samples)); append_define(source, "XE_TRANSFER_SOURCE_TEXTURE_INDEX", 0); append_define(source, "XE_TRANSFER_STENCIL_TEXTURE_INDEX", 1); append_define(source, "XE_TRANSFER_HOST_DEPTH_TEXTURE_INDEX", @@ -5410,6 +6336,9 @@ MTL::RenderPipelineState* MetalRenderTargetCache::GetOrCreateTransferPipelines( output == TransferOutput::kDepth ? 1 : 0); append_define(source, "XE_TRANSFER_OUTPUT_STENCIL_BIT", output == TransferOutput::kStencilBit ? 1 : 0); + append_define(source, "XE_TRANSFER_TILE_INSTANCED", tile_instanced ? 1 : 0); + append_define(source, "XE_TRANSFER_FAST_DIVMOD", + ::cvars::metal_transfer_fast_divmod ? 1 : 0); append_define(source, "XE_FMT_8_8_8_8", uint32_t(xenos::ColorRenderTargetFormat::k_8_8_8_8)); append_define(source, "XE_FMT_8_8_8_8_GAMMA", @@ -5473,11 +6402,27 @@ struct TransferShaderConstants { uint msaa_2x_supported; uint tile_width_samples; uint tile_height_samples; + uint dest_tile_width_pixels; + uint dest_tile_height_pixels; + float dest_tile_width_pixels_inv; + float dest_tile_height_pixels_inv; + float source_pitch_tiles_inv; + float host_depth_source_pitch_tiles_inv; + float dest_pixel_to_ndc_x; + float dest_pixel_to_ndc_y; uint dest_sample_id; uint stencil_mask; uint stencil_clear; }; +struct TransferTileInstance { + float2 tile_origin; + uint tile_index; + uint padding; + uint2 source_base; + uint2 host_base; +}; + constant uint kEdramTileCount = 2048u; inline uint XeBitFieldMask(uint count) { @@ -5545,6 +6490,19 @@ inline float XePWLGammaToLinear(float value) { return linear * (1.0f / 1023.0f); } +inline void XeFastDivMod(uint x, uint w, float inv_w, thread uint& q, + thread uint& r) { + q = uint(float(x) * inv_w); + r = x - q * w; + if (r >= w) { + r -= w; + q += 1u; + } else if (r > x) { + r += w; + q -= 1u; + } +} + inline float XeLinearToPWLGamma(float value) { float clamped = XeSaturateNoNaN(value); float scale; @@ -5677,12 +6635,43 @@ uint XePackColorRGB10A2Float(float4 color) { struct VSOut { float4 position [[position]]; + float2 tile_origin [[flat]]; + uint tile_index [[flat]]; + uint2 source_base [[flat]]; + uint2 host_base [[flat]]; }; vertex VSOut transfer_vs(uint vid [[vertex_id]]) { float2 pt = float2((vid << 1) & 2, vid & 2); VSOut out; out.position = float4(pt * 2.0f - 1.0f, 0.0f, 1.0f); + out.tile_origin = float2(0.0f); + out.tile_index = 0u; + out.source_base = uint2(0u); + out.host_base = uint2(0u); + return out; +} + +vertex VSOut transfer_tile_vs(uint vid [[vertex_id]], + uint iid [[instance_id]], + constant TransferShaderConstants& constants + [[buffer(0)]], + device const TransferTileInstance* instances + [[buffer(1)]]) { + float2 quad = float2(float(vid & 1), float(vid >> 1)); + TransferTileInstance inst = instances[iid]; + float2 tile_size = float2(constants.dest_tile_width_pixels, + constants.dest_tile_height_pixels); + float2 pos_pixel = inst.tile_origin + quad * tile_size; + float2 ndc; + ndc.x = pos_pixel.x * constants.dest_pixel_to_ndc_x - 1.0f; + ndc.y = 1.0f - pos_pixel.y * constants.dest_pixel_to_ndc_y; + VSOut out; + out.position = float4(ndc, 0.0f, 1.0f); + out.tile_origin = inst.tile_origin; + out.tile_index = inst.tile_index; + out.source_base = inst.source_base; + out.host_base = inst.host_base; return out; } @@ -5709,18 +6698,30 @@ vertex VSOut transfer_vs(uint vid [[vertex_id]]) { #endif #endif #else - #if XE_TRANSFER_SOURCE_IS_MULTISAMPLE - #define XE_TRANSFER_SOURCE_PARAMS \ - , texture2d_ms xe_transfer_source_depth \ - [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]], \ - texture2d_ms xe_transfer_source_stencil \ - [[texture(XE_TRANSFER_STENCIL_TEXTURE_INDEX)]] + #if XE_TRANSFER_SOURCE_NEEDS_STENCIL + #if XE_TRANSFER_SOURCE_IS_MULTISAMPLE + #define XE_TRANSFER_SOURCE_PARAMS \ + , texture2d_ms xe_transfer_source_depth \ + [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]], \ + texture2d_ms xe_transfer_source_stencil \ + [[texture(XE_TRANSFER_STENCIL_TEXTURE_INDEX)]] + #else + #define XE_TRANSFER_SOURCE_PARAMS \ + , texture2d xe_transfer_source_depth \ + [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]], \ + texture2d xe_transfer_source_stencil \ + [[texture(XE_TRANSFER_STENCIL_TEXTURE_INDEX)]] + #endif #else - #define XE_TRANSFER_SOURCE_PARAMS \ - , texture2d xe_transfer_source_depth \ - [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]], \ - texture2d xe_transfer_source_stencil \ - [[texture(XE_TRANSFER_STENCIL_TEXTURE_INDEX)]] + #if XE_TRANSFER_SOURCE_IS_MULTISAMPLE + #define XE_TRANSFER_SOURCE_PARAMS \ + , texture2d_ms xe_transfer_source_depth \ + [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]] + #else + #define XE_TRANSFER_SOURCE_PARAMS \ + , texture2d xe_transfer_source_depth \ + [[texture(XE_TRANSFER_SOURCE_TEXTURE_INDEX)]] + #endif #endif #endif @@ -5745,7 +6746,11 @@ vertex VSOut transfer_vs(uint vid [[vertex_id]]) { #define XE_TRANSFER_HOST_DEPTH_TEXTURE_PARAM #endif -#define XE_TRANSFER_SAMPLE_ID_PARAM +#if XE_TRANSFER_DEST_IS_MULTISAMPLE && XE_TRANSFER_DEST_SAMPLE_ID_FROM_SAMPLE + #define XE_TRANSFER_SAMPLE_ID_PARAM , uint xe_sample_id [[sample_id]] +#else + #define XE_TRANSFER_SAMPLE_ID_PARAM +#endif #if XE_TRANSFER_OUTPUT_COLOR #if XE_TRANSFER_DEST_IS_UINT @@ -5783,26 +6788,47 @@ fragment TransferColorOut transfer_ps( uint2 dest_pixel = uint2(in.position.xy); uint dest_sample_id = 0u; #if XE_TRANSFER_DEST_IS_MULTISAMPLE - dest_sample_id = constants.dest_sample_id; + #if XE_TRANSFER_DEST_SAMPLE_ID_FROM_SAMPLE + dest_sample_id = xe_sample_id; + #else + dest_sample_id = constants.dest_sample_id; + #endif #endif uint tile_width_samples = constants.tile_width_samples; uint tile_height_samples = constants.tile_height_samples; - uint dest_tile_width_pixels = - tile_width_samples >> - ((constants.dest_is_64bpp != 0u) + - (constants.dest_msaa_samples >= 4u ? 1u : 0u)); - uint dest_tile_height_pixels = - tile_height_samples >> (constants.dest_msaa_samples >= 2u ? 1u : 0u); + uint dest_tile_width_pixels = constants.dest_tile_width_pixels; + uint dest_tile_height_pixels = constants.dest_tile_height_pixels; - uint dest_tile_index_x = dest_pixel.x / dest_tile_width_pixels; - uint dest_tile_pixel_x = dest_pixel.x % dest_tile_width_pixels; - uint dest_tile_index_y = dest_pixel.y / dest_tile_height_pixels; - uint dest_tile_pixel_y = dest_pixel.y % dest_tile_height_pixels; + uint dest_tile_pixel_x = 0u; + uint dest_tile_pixel_y = 0u; + uint dest_tile_index = 0u; +#if XE_TRANSFER_TILE_INSTANCED + uint2 tile_origin = uint2(in.tile_origin); + dest_tile_pixel_x = dest_pixel.x - tile_origin.x; + dest_tile_pixel_y = dest_pixel.y - tile_origin.y; + dest_tile_index = in.tile_index; +#else + uint dest_tile_index_x = 0u; + uint dest_tile_index_y = 0u; +#if XE_TRANSFER_FAST_DIVMOD + XeFastDivMod(dest_pixel.x, dest_tile_width_pixels, + constants.dest_tile_width_pixels_inv, dest_tile_index_x, + dest_tile_pixel_x); + XeFastDivMod(dest_pixel.y, dest_tile_height_pixels, + constants.dest_tile_height_pixels_inv, dest_tile_index_y, + dest_tile_pixel_y); +#else + dest_tile_index_x = dest_pixel.x / dest_tile_width_pixels; + dest_tile_pixel_x = dest_pixel.x % dest_tile_width_pixels; + dest_tile_index_y = dest_pixel.y / dest_tile_height_pixels; + dest_tile_pixel_y = dest_pixel.y % dest_tile_height_pixels; +#endif - uint dest_tile_index = + dest_tile_index = dest_tile_index_x + dest_tile_index_y * constants.address.dest_pitch; +#endif uint source_sample_id = dest_sample_id; uint source_tile_pixel_x = dest_tile_pixel_x; @@ -5810,10 +6836,10 @@ fragment TransferColorOut transfer_ps( uint source_color_half = 0u; bool source_color_half_valid = false; - bool source_is_64bpp = constants.source_is_64bpp != 0u; - bool dest_is_64bpp = constants.dest_is_64bpp != 0u; - uint source_msaa = constants.source_msaa_samples; - uint dest_msaa = constants.dest_msaa_samples; + bool source_is_64bpp = XE_TRANSFER_SOURCE_IS_64BPP != 0u; + bool dest_is_64bpp = XE_TRANSFER_DEST_IS_64BPP != 0u; + uint source_msaa = XE_TRANSFER_SOURCE_MSAA_SAMPLES; + uint dest_msaa = XE_TRANSFER_DEST_MSAA_SAMPLES; bool msaa_2x_supported = constants.msaa_2x_supported != 0u; if (!source_is_64bpp && dest_is_64bpp) { @@ -5937,7 +6963,7 @@ fragment TransferColorOut transfer_ps( uint source_pixel_width_dwords_log2 = (source_msaa >= 4u ? 1u : 0u) + (source_is_64bpp ? 1u : 0u); - if ((constants.source_is_depth != 0u) != (constants.dest_is_depth != 0u)) { + if ((XE_TRANSFER_SOURCE_IS_DEPTH != 0u) != (XE_TRANSFER_DEST_IS_DEPTH != 0u)) { uint source_32bpp_tile_half_pixels = tile_width_samples >> (1u + source_pixel_width_dwords_log2); if (source_tile_pixel_x < source_32bpp_tile_half_pixels) { @@ -5947,21 +6973,30 @@ fragment TransferColorOut transfer_ps( } } + uint source_pixel_x = 0u; + uint source_pixel_y = 0u; +#if XE_TRANSFER_TILE_INSTANCED + source_pixel_x = in.source_base.x + source_tile_pixel_x; + source_pixel_y = in.source_base.y + source_tile_pixel_y; +#else uint source_tile_index = uint(int(dest_tile_index) + constants.address.source_to_dest) & (kEdramTileCount - 1u); uint source_pitch_tiles = constants.address.source_pitch; - uint source_tile_index_y = source_tile_index / source_pitch_tiles; - uint source_tile_index_x = source_tile_index % source_pitch_tiles; - - uint source_pixel_x = + uint source_tile_index_y = 0u; + uint source_tile_index_x = 0u; + XeFastDivMod(source_tile_index, source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_index_y, + source_tile_index_x); + source_pixel_x = source_tile_index_x * (tile_width_samples >> source_pixel_width_dwords_log2) + source_tile_pixel_x; - uint source_pixel_y = + source_pixel_y = source_tile_index_y * (tile_height_samples >> (source_msaa >= 2u ? 1u : 0u)) + source_tile_pixel_y; +#endif bool load_two = !source_is_64bpp && dest_is_64bpp; uint source_pixel_x1 = source_pixel_x; @@ -6030,6 +7065,7 @@ fragment TransferColorOut transfer_ps( uint2(source_pixel_x1, source_pixel_y)).r; #endif } +#if XE_TRANSFER_SOURCE_NEEDS_STENCIL uint source_stencil0 = #if XE_TRANSFER_SOURCE_IS_MULTISAMPLE xe_transfer_source_stencil.read(uint2(source_pixel_x, source_pixel_y), @@ -6047,12 +7083,16 @@ fragment TransferColorOut transfer_ps( uint2(source_pixel_x1, source_pixel_y)).r; #endif } +#else + uint source_stencil0 = 0u; + uint source_stencil1 = 0u; +#endif #endif #if XE_TRANSFER_SOURCE_IS_COLOR if (source_is_64bpp && !dest_is_64bpp && source_color_half_valid) { uint source_component_count = 0u; - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_32_FLOAT: source_component_count = 1u; break; @@ -6095,7 +7135,7 @@ fragment TransferColorOut transfer_ps( uint2 packed64 = uint2(0u); #if XE_TRANSFER_SOURCE_IS_COLOR #if XE_TRANSFER_SOURCE_IS_UINT - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_16_16: case XE_FMT_16_16_FLOAT: packed64.x = source_color0[0] | (source_color0[1] << 16u); @@ -6120,7 +7160,7 @@ fragment TransferColorOut transfer_ps( break; } #else - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_8_8_8_8_GAMMA: { #if XE_GAMMA_RT_AS_UNORM16 float4 gamma_color0 = source_color0; @@ -6165,7 +7205,7 @@ fragment TransferColorOut transfer_ps( #else uint depth24_0 = 0u; uint depth24_1 = 0u; - if (constants.source_format == XE_FMT_D24FS8) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_D24FS8) { bool round_depth = constants.depth_round != 0u; depth24_0 = XeFloat32To20e4(source_depth0 * 2.0f, round_depth); depth24_1 = XeFloat32To20e4(source_depth1 * 2.0f, round_depth); @@ -6179,7 +7219,7 @@ fragment TransferColorOut transfer_ps( packed64.y = (depth24_1 << 8u) | (source_stencil1 & 0xFFu); #endif - if (constants.dest_format == XE_FMT_32_32_FLOAT) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_32_32_FLOAT) { #if XE_TRANSFER_DEST_IS_UINT out_color = uint4(packed64.x, packed64.y, 0u, 0u); #else @@ -6200,11 +7240,11 @@ fragment TransferColorOut transfer_ps( uint packed32 = 0u; #if XE_TRANSFER_SOURCE_IS_COLOR #if XE_TRANSFER_SOURCE_IS_UINT - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_16_16: case XE_FMT_16_16_FLOAT: - if (constants.dest_format == XE_FMT_16_16 || - constants.dest_format == XE_FMT_16_16_FLOAT) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_16_16 || + XE_TRANSFER_DEST_FORMAT == XE_FMT_16_16_FLOAT) { #if XE_TRANSFER_DEST_IS_UINT out_color = uint4(source_color0[0], source_color0[1], 0u, 0u); #else @@ -6218,8 +7258,8 @@ fragment TransferColorOut transfer_ps( break; case XE_FMT_16_16_16_16: case XE_FMT_16_16_16_16_FLOAT: - if (constants.dest_format == XE_FMT_16_16 || - constants.dest_format == XE_FMT_16_16_FLOAT) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_16_16 || + XE_TRANSFER_DEST_FORMAT == XE_FMT_16_16_FLOAT) { #if XE_TRANSFER_DEST_IS_UINT out_color = uint4(source_color0[0], source_color0[1], 0u, 0u); #else @@ -6240,16 +7280,16 @@ fragment TransferColorOut transfer_ps( break; } #else - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_8_8_8_8: case XE_FMT_8_8_8_8_GAMMA: { float4 color = source_color0; #if XE_GAMMA_RT_AS_UNORM16 - if (constants.source_format == XE_FMT_8_8_8_8_GAMMA && - (constants.dest_format == XE_FMT_8_8_8_8 || - constants.dest_format == XE_FMT_8_8_8_8_GAMMA) && - constants.dest_format != constants.source_format) { - if (constants.dest_format == XE_FMT_8_8_8_8) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_8_8_8_8_GAMMA && + (XE_TRANSFER_DEST_FORMAT == XE_FMT_8_8_8_8 || + XE_TRANSFER_DEST_FORMAT == XE_FMT_8_8_8_8_GAMMA) && + XE_TRANSFER_DEST_FORMAT != XE_TRANSFER_SOURCE_FORMAT) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_8_8_8_8) { color.rgb = XeLinearToPWLGamma3(color.rgb); } else { color.rgb = XePWLGammaToLinear3(color.rgb); @@ -6257,25 +7297,25 @@ fragment TransferColorOut transfer_ps( } #endif #if !XE_TRANSFER_DEST_IS_UINT - if (constants.dest_format == XE_FMT_8_8_8_8 || - constants.dest_format == XE_FMT_8_8_8_8_GAMMA) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_8_8_8_8 || + XE_TRANSFER_DEST_FORMAT == XE_FMT_8_8_8_8_GAMMA) { out_color = color; wrote_direct = true; } else #endif { #if XE_GAMMA_RT_AS_UNORM16 - if (constants.source_format == XE_FMT_8_8_8_8_GAMMA) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_8_8_8_8_GAMMA) { color.rgb = XeLinearToPWLGamma3(color.rgb); } #endif uint packed_component_offset = 0u; - if (constants.dest_is_depth != 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH != 0u) { packed_component_offset = 1u; } packed32 = XePackUnorm(color[packed_component_offset], 255.0f); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed32 |= XePackUnorm(color[packed_component_offset + 1], 255.0f) << 8u; packed32 |= XePackUnorm(color[packed_component_offset + 2], @@ -6288,15 +7328,15 @@ fragment TransferColorOut transfer_ps( case XE_FMT_2_10_10_10: case XE_FMT_2_10_10_10_AS_10_10_10_10: #if !XE_TRANSFER_DEST_IS_UINT - if (constants.dest_format == XE_FMT_2_10_10_10 || - constants.dest_format == XE_FMT_2_10_10_10_AS_10_10_10_10) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_2_10_10_10 || + XE_TRANSFER_DEST_FORMAT == XE_FMT_2_10_10_10_AS_10_10_10_10) { out_color = source_color0; wrote_direct = true; } else #endif { packed32 = XePackUnorm(source_color0[0], 1023.0f); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed32 |= XePackUnorm(source_color0[1], 1023.0f) << 10u; packed32 |= XePackUnorm(source_color0[2], 1023.0f) << 20u; packed32 |= XePackUnorm(source_color0[3], 3.0f) << 30u; @@ -6306,15 +7346,15 @@ fragment TransferColorOut transfer_ps( case XE_FMT_2_10_10_10_FLOAT: case XE_FMT_2_10_10_10_FLOAT_AS_16_16_16_16: #if !XE_TRANSFER_DEST_IS_UINT - if (constants.dest_format == XE_FMT_2_10_10_10_FLOAT || - constants.dest_format == XE_FMT_2_10_10_10_FLOAT_AS_16_16_16_16) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_2_10_10_10_FLOAT || + XE_TRANSFER_DEST_FORMAT == XE_FMT_2_10_10_10_FLOAT_AS_16_16_16_16) { out_color = source_color0; wrote_direct = true; } else #endif { packed32 = XeUnclampedFloat32To7e3(source_color0[0]); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed32 |= XeUnclampedFloat32To7e3(source_color0[1]) << 10u; packed32 |= XeUnclampedFloat32To7e3(source_color0[2]) << 20u; packed32 |= XePackUnorm(source_color0[3], 3.0f) << 30u; @@ -6331,8 +7371,8 @@ fragment TransferColorOut transfer_ps( } #endif #else - if (constants.dest_is_depth != 0u && - constants.dest_format == constants.source_format) { + if (XE_TRANSFER_DEST_IS_DEPTH != 0u && + XE_TRANSFER_DEST_FORMAT == XE_TRANSFER_SOURCE_FORMAT) { TransferColorOut out; out.color = XeTransferColorOutType(source_depth0); #if XE_TRANSFER_DEST_IS_MULTISAMPLE @@ -6340,20 +7380,20 @@ fragment TransferColorOut transfer_ps( #endif return out; } - if (constants.source_format == XE_FMT_D24FS8) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_D24FS8) { bool round_depth = constants.depth_round != 0u; packed32 = XeFloat32To20e4(source_depth0 * 2.0f, round_depth); } else { packed32 = XeRoundToNearestEven( clamp(source_depth0, 0.0f, 1.0f) * 16777215.0f); } - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed32 = (packed32 << 8u) | (source_stencil0 & 0xFFu); } #endif if (!wrote_direct) { - switch (constants.dest_format) { + switch (XE_TRANSFER_DEST_FORMAT) { case XE_FMT_8_8_8_8: { #if XE_TRANSFER_DEST_IS_UINT out_color = uint4(packed32, 0u, 0u, 0u); @@ -6457,35 +7497,56 @@ fragment TransferDepthOut transfer_ps( uint2 dest_pixel = uint2(in.position.xy); uint dest_sample_id = 0u; #if XE_TRANSFER_DEST_IS_MULTISAMPLE - dest_sample_id = constants.dest_sample_id; + #if XE_TRANSFER_DEST_SAMPLE_ID_FROM_SAMPLE + dest_sample_id = xe_sample_id; + #else + dest_sample_id = constants.dest_sample_id; + #endif #endif uint tile_width_samples = constants.tile_width_samples; uint tile_height_samples = constants.tile_height_samples; - uint dest_tile_width_pixels = - tile_width_samples >> - ((constants.dest_is_64bpp != 0u) + - (constants.dest_msaa_samples >= 4u ? 1u : 0u)); - uint dest_tile_height_pixels = - tile_height_samples >> (constants.dest_msaa_samples >= 2u ? 1u : 0u); + uint dest_tile_width_pixels = constants.dest_tile_width_pixels; + uint dest_tile_height_pixels = constants.dest_tile_height_pixels; - uint dest_tile_index_x = dest_pixel.x / dest_tile_width_pixels; - uint dest_tile_pixel_x = dest_pixel.x % dest_tile_width_pixels; - uint dest_tile_index_y = dest_pixel.y / dest_tile_height_pixels; - uint dest_tile_pixel_y = dest_pixel.y % dest_tile_height_pixels; + uint dest_tile_pixel_x = 0u; + uint dest_tile_pixel_y = 0u; + uint dest_tile_index = 0u; +#if XE_TRANSFER_TILE_INSTANCED + uint2 tile_origin = uint2(in.tile_origin); + dest_tile_pixel_x = dest_pixel.x - tile_origin.x; + dest_tile_pixel_y = dest_pixel.y - tile_origin.y; + dest_tile_index = in.tile_index; +#else + uint dest_tile_index_x = 0u; + uint dest_tile_index_y = 0u; +#if XE_TRANSFER_FAST_DIVMOD + XeFastDivMod(dest_pixel.x, dest_tile_width_pixels, + constants.dest_tile_width_pixels_inv, dest_tile_index_x, + dest_tile_pixel_x); + XeFastDivMod(dest_pixel.y, dest_tile_height_pixels, + constants.dest_tile_height_pixels_inv, dest_tile_index_y, + dest_tile_pixel_y); +#else + dest_tile_index_x = dest_pixel.x / dest_tile_width_pixels; + dest_tile_pixel_x = dest_pixel.x % dest_tile_width_pixels; + dest_tile_index_y = dest_pixel.y / dest_tile_height_pixels; + dest_tile_pixel_y = dest_pixel.y % dest_tile_height_pixels; +#endif - uint dest_tile_index = + dest_tile_index = dest_tile_index_x + dest_tile_index_y * constants.address.dest_pitch; +#endif uint source_sample_id = dest_sample_id; uint source_tile_pixel_x = dest_tile_pixel_x; uint source_tile_pixel_y = dest_tile_pixel_y; - bool source_is_64bpp = constants.source_is_64bpp != 0u; - bool dest_is_64bpp = constants.dest_is_64bpp != 0u; - uint source_msaa = constants.source_msaa_samples; - uint dest_msaa = constants.dest_msaa_samples; + bool source_is_64bpp = XE_TRANSFER_SOURCE_IS_64BPP != 0u; + bool dest_is_64bpp = XE_TRANSFER_DEST_IS_64BPP != 0u; + uint source_msaa = XE_TRANSFER_SOURCE_MSAA_SAMPLES; + uint dest_msaa = XE_TRANSFER_DEST_MSAA_SAMPLES; bool msaa_2x_supported = constants.msaa_2x_supported != 0u; if (!source_is_64bpp && dest_is_64bpp) { @@ -6605,7 +7666,7 @@ fragment TransferDepthOut transfer_ps( uint source_pixel_width_dwords_log2 = (source_msaa >= 4u ? 1u : 0u) + (source_is_64bpp ? 1u : 0u); - if ((constants.source_is_depth != 0u) != (constants.dest_is_depth != 0u)) { + if ((XE_TRANSFER_SOURCE_IS_DEPTH != 0u) != (XE_TRANSFER_DEST_IS_DEPTH != 0u)) { uint source_32bpp_tile_half_pixels = tile_width_samples >> (1u + source_pixel_width_dwords_log2); if (source_tile_pixel_x < source_32bpp_tile_half_pixels) { @@ -6615,21 +7676,30 @@ fragment TransferDepthOut transfer_ps( } } + uint source_pixel_x = 0u; + uint source_pixel_y = 0u; +#if XE_TRANSFER_TILE_INSTANCED + source_pixel_x = in.source_base.x + source_tile_pixel_x; + source_pixel_y = in.source_base.y + source_tile_pixel_y; +#else uint source_tile_index = uint(int(dest_tile_index) + constants.address.source_to_dest) & (kEdramTileCount - 1u); uint source_pitch_tiles = constants.address.source_pitch; - uint source_tile_index_y = source_tile_index / source_pitch_tiles; - uint source_tile_index_x = source_tile_index % source_pitch_tiles; - - uint source_pixel_x = + uint source_tile_index_y = 0u; + uint source_tile_index_x = 0u; + XeFastDivMod(source_tile_index, source_pitch_tiles, + constants.source_pitch_tiles_inv, source_tile_index_y, + source_tile_index_x); + source_pixel_x = source_tile_index_x * (tile_width_samples >> source_pixel_width_dwords_log2) + source_tile_pixel_x; - uint source_pixel_y = + source_pixel_y = source_tile_index_y * (tile_height_samples >> (source_msaa >= 2u ? 1u : 0u)) + source_tile_pixel_y; +#endif bool load_two = !source_is_64bpp && dest_is_64bpp; uint source_pixel_x1 = source_pixel_x; @@ -6698,6 +7768,7 @@ fragment TransferDepthOut transfer_ps( uint2(source_pixel_x1, source_pixel_y)).r; #endif } +#if XE_TRANSFER_SOURCE_NEEDS_STENCIL uint source_stencil0 = #if XE_TRANSFER_SOURCE_IS_MULTISAMPLE xe_transfer_source_stencil.read(uint2(source_pixel_x, source_pixel_y), @@ -6715,6 +7786,10 @@ fragment TransferDepthOut transfer_ps( uint2(source_pixel_x1, source_pixel_y)).r; #endif } +#else + uint source_stencil0 = 0u; + uint source_stencil1 = 0u; +#endif #endif uint packed = 0u; @@ -6723,7 +7798,7 @@ fragment TransferDepthOut transfer_ps( #endif #if XE_TRANSFER_SOURCE_IS_COLOR #if XE_TRANSFER_SOURCE_IS_UINT - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_16_16: case XE_FMT_16_16_FLOAT: case XE_FMT_16_16_16_16: @@ -6739,24 +7814,24 @@ fragment TransferDepthOut transfer_ps( break; } #else - switch (constants.source_format) { + switch (XE_TRANSFER_SOURCE_FORMAT) { case XE_FMT_8_8_8_8: case XE_FMT_8_8_8_8_GAMMA: { float4 color = source_color0; - if (constants.source_format == XE_FMT_8_8_8_8_GAMMA) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_8_8_8_8_GAMMA) { #if XE_GAMMA_RT_AS_UNORM16 || XE_TRANSFER_OUTPUT_STENCIL_BIT color.rgb = XeLinearToPWLGamma3(color.rgb); #endif } uint packed_component_offset = 0u; - if (constants.dest_is_depth != 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH != 0u) { packed_component_offset = 1u; #if !XE_TRANSFER_OUTPUT_STENCIL_BIT packed_only_depth = true; #endif } packed = XePackUnorm(color[packed_component_offset], 255.0f); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed |= XePackUnorm(color[packed_component_offset + 1], 255.0f) << 8u; packed |= XePackUnorm(color[packed_component_offset + 2], @@ -6768,7 +7843,7 @@ fragment TransferDepthOut transfer_ps( case XE_FMT_2_10_10_10: case XE_FMT_2_10_10_10_AS_10_10_10_10: { packed = XePackUnorm(source_color0[0], 1023.0f); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed |= XePackUnorm(source_color0[1], 1023.0f) << 10u; packed |= XePackUnorm(source_color0[2], 1023.0f) << 20u; packed |= XePackUnorm(source_color0[3], 3.0f) << 30u; @@ -6777,7 +7852,7 @@ fragment TransferDepthOut transfer_ps( case XE_FMT_2_10_10_10_FLOAT: case XE_FMT_2_10_10_10_FLOAT_AS_16_16_16_16: { packed = XeUnclampedFloat32To7e3(source_color0[0]); - if (constants.dest_is_depth == 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH == 0u) { packed |= XeUnclampedFloat32To7e3(source_color0[1]) << 10u; packed |= XeUnclampedFloat32To7e3(source_color0[2]) << 20u; packed |= XePackUnorm(source_color0[3], 3.0f) << 30u; @@ -6793,14 +7868,14 @@ fragment TransferDepthOut transfer_ps( } #endif #else - if (constants.source_format == XE_FMT_D24FS8) { + if (XE_TRANSFER_SOURCE_FORMAT == XE_FMT_D24FS8) { bool round_depth = constants.depth_round != 0u; packed = XeFloat32To20e4(source_depth0 * 2.0f, round_depth); } else { packed = XeRoundToNearestEven( clamp(source_depth0, 0.0f, 1.0f) * 16777215.0f); } - if (constants.dest_is_depth != 0u) { + if (XE_TRANSFER_DEST_IS_DEPTH != 0u) { #if !XE_TRANSFER_OUTPUT_STENCIL_BIT packed_only_depth = true; #endif @@ -6808,6 +7883,9 @@ fragment TransferDepthOut transfer_ps( packed = (packed << 8u) | (source_stencil0 & 0xFFu); } #endif +#if XE_TRANSFER_OUTPUT_STENCIL_BIT && !XE_TRANSFER_SOURCE_IS_COLOR + packed = source_stencil0; +#endif #if XE_TRANSFER_OUTPUT_STENCIL_BIT if (constants.stencil_clear == 0u) { @@ -6834,7 +7912,7 @@ fragment TransferDepthOut transfer_ps( uint host_tile_pixel_x = dest_tile_pixel_x; uint host_tile_pixel_y = dest_tile_pixel_y; uint host_sample_id = dest_sample_id; - uint host_msaa = constants.host_depth_source_msaa_samples; + uint host_msaa = XE_TRANSFER_HOST_DEPTH_MSAA_SAMPLES; if (host_msaa != dest_msaa) { if (host_msaa >= 4u) { @@ -6895,20 +7973,30 @@ fragment TransferDepthOut transfer_ps( } } + uint host_pixel_x = 0u; + uint host_pixel_y = 0u; +#if XE_TRANSFER_TILE_INSTANCED + host_pixel_x = in.host_base.x + host_tile_pixel_x; + host_pixel_y = in.host_base.y + host_tile_pixel_y; +#else uint host_tile_index = uint(int(dest_tile_index) + constants.host_depth_address.source_to_dest) & (kEdramTileCount - 1u); uint host_pitch_tiles = constants.host_depth_address.source_pitch; - uint host_tile_index_y = host_tile_index / host_pitch_tiles; - uint host_tile_index_x = host_tile_index % host_pitch_tiles; - uint host_pixel_x = + uint host_tile_index_y = 0u; + uint host_tile_index_x = 0u; + XeFastDivMod(host_tile_index, host_pitch_tiles, + constants.host_depth_source_pitch_tiles_inv, host_tile_index_y, + host_tile_index_x); + host_pixel_x = host_tile_index_x * (tile_width_samples >> (host_msaa >= 4u ? 1u : 0u)) + host_tile_pixel_x; - uint host_pixel_y = + host_pixel_y = host_tile_index_y * (tile_height_samples >> (host_msaa >= 2u ? 1u : 0u)) + host_tile_pixel_y; +#endif #if XE_TRANSFER_HOST_DEPTH_IS_MULTISAMPLE host_depth32 = xe_transfer_host_depth.read( @@ -6945,7 +8033,7 @@ fragment TransferDepthOut transfer_ps( #endif float fragment_depth = 0.0f; - if (constants.dest_format == XE_FMT_D24FS8) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_D24FS8) { float guest_depth32 = XeFloat20e4To32(guest_depth24, true); fragment_depth = guest_depth32; } else { @@ -6954,7 +8042,7 @@ fragment TransferDepthOut transfer_ps( if (has_host_depth) { uint host_depth24 = 0u; - if (constants.dest_format == XE_FMT_D24FS8) { + if (XE_TRANSFER_DEST_FORMAT == XE_FMT_D24FS8) { bool round_depth = constants.depth_round != 0u; host_depth24 = XeFloat32To20e4(host_depth32, round_depth); } else { @@ -6994,7 +8082,9 @@ fragment TransferDepthOut transfer_ps( return nullptr; } - auto vs_name = NS::String::string("transfer_vs", NS::UTF8StringEncoding); + auto vs_name = + NS::String::string(tile_instanced ? "transfer_tile_vs" : "transfer_vs", + NS::UTF8StringEncoding); auto ps_name = NS::String::string("transfer_ps", NS::UTF8StringEncoding); MTL::Function* vs = lib->newFunction(vs_name); MTL::Function* ps = lib->newFunction(ps_name); @@ -7049,7 +8139,7 @@ fragment TransferDepthOut transfer_ps( return nullptr; } - transfer_pipelines_.emplace(key, pipeline); + pipeline_map.emplace(key, pipeline); return pipeline; } diff --git a/src/xenia/gpu/metal/metal_render_target_cache.h b/src/xenia/gpu/metal/metal_render_target_cache.h index 47a0e78ed..cfba15027 100644 --- a/src/xenia/gpu/metal/metal_render_target_cache.h +++ b/src/xenia/gpu/metal/metal_render_target_cache.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -10,6 +10,7 @@ #ifndef XENIA_GPU_METAL_METAL_RENDER_TARGET_CACHE_H_ #define XENIA_GPU_METAL_METAL_RENDER_TARGET_CACHE_H_ +#include #include #include #include @@ -53,13 +54,23 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { MTL::Texture* msaa_transfer_texture() const { return msaa_transfer_texture_ ? msaa_transfer_texture_ : msaa_texture_; } + MTL::Texture* stencil_view() const { return stencil_view_; } + void SetStencilView(MTL::Texture* view) { stencil_view_ = view; } void SetTemporarySortIndex(uint32_t index) { temporary_sort_index_ = index; } uint32_t temporary_sort_index() const { return temporary_sort_index_; } - void SetTexture(MTL::Texture* texture) { texture_ = texture; } + void SetTexture(MTL::Texture* texture) { + if (texture_ != texture) { + if (stencil_view_) { + stencil_view_->release(); + stencil_view_ = nullptr; + } + texture_ = texture; + } + } void SetMsaaTexture(MTL::Texture* texture) { msaa_texture_ = texture; } void SetDrawTexture(MTL::Texture* texture) { draw_texture_ = texture; } void SetTransferTexture(MTL::Texture* texture) { @@ -86,6 +97,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { MTL::Texture* transfer_texture_ = nullptr; MTL::Texture* msaa_draw_texture_ = nullptr; MTL::Texture* msaa_transfer_texture_ = nullptr; + MTL::Texture* stencil_view_ = nullptr; uint32_t temporary_sort_index_ = UINT32_MAX; bool needs_initial_clear_ = true; }; @@ -223,7 +235,6 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { MTL::ComputePipelineState* edram_dump_color_32bpp_1xmsaa_pipeline_ = nullptr; MTL::ComputePipelineState* edram_dump_color_32bpp_2xmsaa_pipeline_ = nullptr; MTL::ComputePipelineState* edram_dump_color_32bpp_4xmsaa_pipeline_ = nullptr; - // EDRAM blend compute shaders (host RT -> EDRAM with blend/keep mask). // Color, 64bpp. MTL::ComputePipelineState* edram_dump_color_64bpp_1xmsaa_pipeline_ = nullptr; MTL::ComputePipelineState* edram_dump_color_64bpp_2xmsaa_pipeline_ = nullptr; @@ -260,8 +271,6 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { // Host depth store compute shaders (1x/2x/4x MSAA). MTL::ComputePipelineState* host_depth_store_pipelines_[3] = {}; - // Per-draw ordered-blend coverage attachment. - // Transfer shaders (host RT ownership transfers) - modeled after D3D12. // TransferMode list mirrors D3D12RenderTargetCache::TransferMode so logs and @@ -284,6 +293,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { xenos::MsaaSamples host_depth_source_msaa_samples; uint32_t source_resource_format; uint32_t dest_resource_format; + uint32_t dest_sample_id_from_sample; uint32_t host_depth_source_is_copy; bool operator==(const TransferShaderKey& other) const { @@ -294,6 +304,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { other.host_depth_source_msaa_samples && source_resource_format == other.source_resource_format && dest_resource_format == other.dest_resource_format && + dest_sample_id_from_sample == other.dest_sample_id_from_sample && host_depth_source_is_copy == other.host_depth_source_is_copy; } bool operator!=(const TransferShaderKey& other) const { @@ -320,6 +331,9 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { if (dest_resource_format != other.dest_resource_format) { return dest_resource_format < other.dest_resource_format; } + if (dest_sample_id_from_sample != other.dest_sample_id_from_sample) { + return dest_sample_id_from_sample < other.dest_sample_id_from_sample; + } return host_depth_source_is_copy < other.host_depth_source_is_copy; } @@ -331,7 +345,8 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { h ^= (size_t(key.host_depth_source_msaa_samples) << 12); h ^= (size_t(key.source_resource_format) << 16); h ^= (size_t(key.dest_resource_format) << 24); - h ^= (size_t(key.host_depth_source_is_copy) << 28); + h ^= (size_t(key.dest_sample_id_from_sample) << 28); + h ^= (size_t(key.host_depth_source_is_copy) << 29); return h ^ (h >> 16); } }; @@ -369,10 +384,22 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { std::unordered_map transfer_pipelines_; + std::unordered_map + transfer_tile_pipelines_; std::vector transfer_invocations_; MTL::Library* transfer_library_ = nullptr; std::unordered_map transfer_clear_pipelines_; + static constexpr uint32_t kTransferInstanceBufferCount = 3; + std::array + transfer_tile_instance_buffers_ = {}; + std::array + transfer_tile_instance_buffer_sizes_ = {}; + std::array, kTransferInstanceBufferCount> + transfer_tile_instance_retired_buffers_ = {}; + uint64_t transfer_tile_instance_buffer_frame_id_ = 0; + size_t transfer_tile_instance_buffer_offset_ = 0; MTL::DepthStencilState* transfer_depth_state_ = nullptr; MTL::DepthStencilState* transfer_depth_state_none_ = nullptr; MTL::DepthStencilState* transfer_depth_clear_state_ = nullptr; @@ -400,6 +427,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { // Render pass descriptor cache MTL::RenderPassDescriptor* cached_render_pass_descriptor_ = nullptr; bool render_pass_descriptor_dirty_ = true; + uint32_t cached_render_pass_descriptor_sample_count_ = 0; // Dummy render target for when no render targets are bound struct DummyColorTargetEntry { @@ -423,6 +451,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { MTL::Texture* CreateDepthTexture(uint32_t width, uint32_t height, xenos::DepthRenderTargetFormat format, uint32_t samples); + MTL::Texture* GetStencilTextureView(MetalRenderTarget* render_target); MTL::PixelFormat GetColorResourcePixelFormat( xenos::ColorRenderTargetFormat format) const; @@ -441,7 +470,7 @@ class MetalRenderTargetCache final : public gpu::RenderTargetCache { // D3D12RenderTargetCache::GetOrCreateTransferPipelines. MTL::RenderPipelineState* GetOrCreateTransferPipelines( const TransferShaderKey& key, MTL::PixelFormat dest_format, - bool dest_is_uint); + bool dest_is_uint, bool tile_instanced); MTL::RenderPipelineState* GetOrCreateTransferClearPipeline( MTL::PixelFormat dest_format, bool dest_is_uint, bool is_depth, uint32_t sample_count); diff --git a/src/xenia/gpu/metal/metal_shader.cc b/src/xenia/gpu/metal/metal_shader.cc index 3cb9d77b2..1a28317ef 100644 --- a/src/xenia/gpu/metal/metal_shader.cc +++ b/src/xenia/gpu/metal/metal_shader.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_shader.h b/src/xenia/gpu/metal/metal_shader.h index 261e8f5eb..cd02988f0 100644 --- a/src/xenia/gpu/metal/metal_shader.h +++ b/src/xenia/gpu/metal/metal_shader.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_shader_cache.cc b/src/xenia/gpu/metal/metal_shader_cache.cc index 11033cd74..a11f4f0a8 100644 --- a/src/xenia/gpu/metal/metal_shader_cache.cc +++ b/src/xenia/gpu/metal/metal_shader_cache.cc @@ -1,7 +1,10 @@ /** - ******************************************************************************* + ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * - ******************************************************************************* + ****************************************************************************** + * Copyright 2026 Ben Vanik. All rights reserved. * + * Released under the BSD license - see LICENSE in the root for more details. * + ****************************************************************************** */ #include "xenia/gpu/metal/metal_shader_cache.h" diff --git a/src/xenia/gpu/metal/metal_shader_cache.h b/src/xenia/gpu/metal/metal_shader_cache.h index a1ac44d8e..308e0d7e7 100644 --- a/src/xenia/gpu/metal/metal_shader_cache.h +++ b/src/xenia/gpu/metal/metal_shader_cache.h @@ -1,7 +1,10 @@ /** - ******************************************************************************* + ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * - ******************************************************************************* + ****************************************************************************** + * Copyright 2026 Ben Vanik. All rights reserved. * + * Released under the BSD license - see LICENSE in the root for more details. * + ****************************************************************************** */ #ifndef XENIA_GPU_METAL_METAL_SHADER_CACHE_H_ diff --git a/src/xenia/gpu/metal/metal_shader_converter.cc b/src/xenia/gpu/metal/metal_shader_converter.cc index 14e850364..09a9c92ef 100644 --- a/src/xenia/gpu/metal/metal_shader_converter.cc +++ b/src/xenia/gpu/metal/metal_shader_converter.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_shader_converter.h b/src/xenia/gpu/metal/metal_shader_converter.h index 2dc7bf912..1e809d864 100644 --- a/src/xenia/gpu/metal/metal_shader_converter.h +++ b/src/xenia/gpu/metal/metal_shader_converter.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_shared_memory.cc b/src/xenia/gpu/metal/metal_shared_memory.cc index 84e7d8b8a..5d609fc21 100644 --- a/src/xenia/gpu/metal/metal_shared_memory.cc +++ b/src/xenia/gpu/metal/metal_shared_memory.cc @@ -1,3 +1,12 @@ +/** + ****************************************************************************** + * Xenia : Xbox 360 Emulator Research Project * + ****************************************************************************** + * Copyright 2026 Ben Vanik. All rights reserved. * + * Released under the BSD license - see LICENSE in the root for more details. * + ****************************************************************************** + */ + #include "xenia/gpu/metal/metal_shared_memory.h" #include "xenia/base/logging.h" diff --git a/src/xenia/gpu/metal/metal_shared_memory.h b/src/xenia/gpu/metal/metal_shared_memory.h index e2b0f218b..114dc1536 100644 --- a/src/xenia/gpu/metal/metal_shared_memory.h +++ b/src/xenia/gpu/metal/metal_shared_memory.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_texture_cache.cc b/src/xenia/gpu/metal/metal_texture_cache.cc index 9c851113b..1144e77e1 100644 --- a/src/xenia/gpu/metal/metal_texture_cache.cc +++ b/src/xenia/gpu/metal/metal_texture_cache.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -93,9 +93,6 @@ #include "xenia/gpu/texture_util.h" #include "xenia/gpu/xenos.h" -DEFINE_bool(metal_texture_dump_png, false, - "Dump some loaded Metal textures as PNG to scratch/gpu (debug).", - "GPU"); DEFINE_bool(metal_force_bc_decompress, false, "Force BC1/2/3/5/DXN decompression to RGBA8/RG8 (debug).", "GPU"); @@ -1971,55 +1968,6 @@ bool MetalTextureCache::UpdateTextureCube(MTL::Texture* texture, return true; } -bool MetalTextureCache::ConvertTextureData(const void* src_data, void* dst_data, - uint32_t width, uint32_t height, - xenos::TextureFormat src_format, - MTL::PixelFormat dst_format) { - // TODO: Implement proper Xbox 360 texture format conversion - // For now, just copy data directly - if (src_data && dst_data) { - uint32_t size = width * height * 4; // Assume 4 bytes per pixel - std::memcpy(dst_data, src_data, size); - return true; - } - return false; -} - -MTL::Texture* MetalTextureCache::CreateDebugTexture(uint32_t width, - uint32_t height) { - MTL::TextureSwizzleChannels rgba = { - MTL::TextureSwizzleRed, MTL::TextureSwizzleGreen, MTL::TextureSwizzleBlue, - MTL::TextureSwizzleAlpha}; - MTL::Texture* texture = - CreateTexture2D(width, height, 1, MTL::PixelFormatRGBA8Unorm, rgba); - if (!texture) { - XELOGE("Failed to create debug texture"); - return nullptr; - } - - // Create checkerboard pattern data - std::vector pixels(width * height); - uint32_t checker_size = 32; // Size of each checker square - - for (uint32_t y = 0; y < height; y++) { - for (uint32_t x = 0; x < width; x++) { - bool is_white = ((x / checker_size) + (y / checker_size)) % 2 == 0; - // Use green/purple pattern to distinguish from pink error color - if (is_white) { - pixels[y * width + x] = 0xFF00FF00; // Green (RGBA) - } else { - pixels[y * width + x] = 0xFF8000FF; // Purple (RGBA) - } - } - } - - // Upload texture data - MTL::Region region = MTL::Region::Make2D(0, 0, width, height); - texture->replaceRegion(region, 0, pixels.data(), width * 4); - - return texture; -} - MTL::Texture* MetalTextureCache::CreateNullTexture2D() { SCOPE_profile_cpu_f("gpu"); @@ -2781,17 +2729,21 @@ bool MetalTextureCache::EnsureScaledResolveBufferRange(uint64_t start_scaled, NS::String::string("XeniaScaledResolveBuffer", NS::UTF8StringEncoding)); if (!overlap_indices.empty()) { - MTL::CommandQueue* queue = command_processor_->GetMetalCommandQueue(); - if (!queue) { - new_buffer->release(); - return false; - } - ScopedAutoreleasePool autorelease_pool; - MTL::CommandBuffer* cmd = queue->commandBuffer(); + MTL::CommandBuffer* cmd = command_processor_->GetCurrentCommandBuffer(); if (!cmd) { - new_buffer->release(); - return false; + MTL::CommandQueue* queue = command_processor_->GetMetalCommandQueue(); + if (!queue) { + new_buffer->release(); + return false; + } + ScopedAutoreleasePool autorelease_pool; + cmd = queue->commandBuffer(); + if (!cmd) { + new_buffer->release(); + return false; + } } + MTL::BlitCommandEncoder* blit = cmd->blitCommandEncoder(); if (!blit) { new_buffer->release(); @@ -2809,8 +2761,10 @@ bool MetalTextureCache::EnsureScaledResolveBufferRange(uint64_t start_scaled, } blit->endEncoding(); - cmd->commit(); - cmd->waitUntilCompleted(); + if (cmd != command_processor_->GetCurrentCommandBuffer()) { + cmd->commit(); + cmd->waitUntilCompleted(); + } } std::vector new_buffers; diff --git a/src/xenia/gpu/metal/metal_texture_cache.h b/src/xenia/gpu/metal/metal_texture_cache.h index 0c461461d..37224f17b 100644 --- a/src/xenia/gpu/metal/metal_texture_cache.h +++ b/src/xenia/gpu/metal/metal_texture_cache.h @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ @@ -60,9 +60,6 @@ class MetalTextureCache : public TextureCache { xenos::TextureFormat format, xenos::Endian endian = xenos::Endian::k8in32); - // Debug texture creation - MTL::Texture* CreateDebugTexture(uint32_t width = 256, uint32_t height = 256); - // Null texture accessors for invalid bindings (following D3D12/Vulkan // pattern) MTL::Texture* GetNullTexture2D() const { return null_texture_2d_; } @@ -119,6 +116,12 @@ class MetalTextureCache : public TextureCache { bool GetCurrentScaledResolveBuffer(MTL::Buffer*& buffer_out, size_t& buffer_offset_out, size_t& buffer_length_out) const; + uint64_t GetCurrentScaledResolveRangeStartScaled() const { + return scaled_resolve_current_range_start_scaled_; + } + uint64_t GetCurrentScaledResolveRangeLengthScaled() const { + return scaled_resolve_current_range_length_scaled_; + } uint32_t GetHostFormatSwizzle(TextureKey key) const override; uint32_t GetMaxHostTextureWidthHeight( xenos::DataDimension dimension) const override; @@ -211,11 +214,6 @@ class MetalTextureCache : public TextureCache { uint64_t length_scaled); void ClearScaledResolveBuffers(); - // Format conversion helpers - bool ConvertTextureData(const void* src_data, void* dst_data, uint32_t width, - uint32_t height, xenos::TextureFormat src_format, - MTL::PixelFormat dst_format); - // Null texture factory methods (following existing CreateTexture pattern) MTL::Texture* CreateNullTexture2D(); MTL::Texture* CreateNullTexture3D(); diff --git a/src/xenia/gpu/metal/metal_trace_dump_main.cc b/src/xenia/gpu/metal/metal_trace_dump_main.cc index 302821ee9..1729d69c3 100644 --- a/src/xenia/gpu/metal/metal_trace_dump_main.cc +++ b/src/xenia/gpu/metal/metal_trace_dump_main.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2024 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */ diff --git a/src/xenia/gpu/metal/metal_trace_viewer_main.cc b/src/xenia/gpu/metal/metal_trace_viewer_main.cc index 02c219372..17c7ecb5a 100644 --- a/src/xenia/gpu/metal/metal_trace_viewer_main.cc +++ b/src/xenia/gpu/metal/metal_trace_viewer_main.cc @@ -2,7 +2,7 @@ ****************************************************************************** * Xenia : Xbox 360 Emulator Research Project * ****************************************************************************** - * Copyright 2025 Ben Vanik. All rights reserved. * + * Copyright 2026 Ben Vanik. All rights reserved. * * Released under the BSD license - see LICENSE in the root for more details. * ****************************************************************************** */