From ce8198288771643c8755d630b21039daca98306e Mon Sep 17 00:00:00 2001 From: Forrest Mark X Date: Sun, 9 Aug 2026 20:41:55 -0500 Subject: [PATCH] Added ability to toggle the async unswizzle system between Off, CPU, GPU Restored older GPU unswizzle system that worked better Added CPU async texture unswizzling for VRAM limited devices Fixed issue with GPU unswizzle that caused the sparse texture to skip slices Attempted to save ram by only allocating space for the slices that have data Improve speed by saving a buffer that doesn't need to be allocated per-sparse texture Attempt to speed up the CPU unswizzle side by batching to reduce loop counts and using SSE2/AVX2 --- src/common/settings.h | 10 +- src/common/settings_enums.h | 1 + src/qt_common/config/shared_translation.cpp | 20 +- src/video_core/CMakeLists.txt | 13 + .../block_linear_unswizzle_3d_bcn.comp | 3 +- .../renderer_opengl/gl_texture_cache.cpp | 6 +- .../renderer_opengl/gl_texture_cache.h | 8 +- .../renderer_vulkan/vk_compute_pass.cpp | 247 ++--- .../renderer_vulkan/vk_compute_pass.h | 9 +- .../renderer_vulkan/vk_texture_cache.cpp | 37 +- .../renderer_vulkan/vk_texture_cache.h | 6 +- .../texture_cache/accelerated_swizzle.h | 78 -- src/video_core/texture_cache/texture_cache.h | 928 +++++++++++++++--- .../texture_cache/texture_cache_base.h | 83 +- src/video_core/textures/decoders.cpp | 133 ++- src/video_core/textures/decoders_avx2.cpp | 157 +++ src/video_core/textures/decoders_avx2.h | 16 + src/video_core/textures/decoders_sse2.cpp | 131 +++ src/video_core/textures/decoders_sse2.h | 16 + 19 files changed, 1434 insertions(+), 468 deletions(-) create mode 100644 src/video_core/textures/decoders_avx2.cpp create mode 100644 src/video_core/textures/decoders_avx2.h create mode 100644 src/video_core/textures/decoders_sse2.cpp create mode 100644 src/video_core/textures/decoders_sse2.h diff --git a/src/common/settings.h b/src/common/settings.h index 6db1cefea9..7e281aa896 100644 --- a/src/common/settings.h +++ b/src/common/settings.h @@ -610,8 +610,14 @@ struct Values { Category::RendererHacks, Specialization::Default}; - SwitchableSetting gpu_unswizzle_enabled{linkage, false, "gpu_unswizzle_enabled", - Category::RendererHacks}; + /*SwitchableSetting gpu_unswizzle_enabled{linkage, false, "gpu_unswizzle_enabled", + Category::RendererHacks};*/ + + SwitchableSetting async_unswizzle_mode{linkage, + AsyncUnswizzleMode::Off, + "async_unswizzle_mode", + Category::RendererHacks, + Specialization::Default}; SwitchableSetting dyna_state{linkage, #if defined(__ANDROID__) diff --git a/src/common/settings_enums.h b/src/common/settings_enums.h index c4895d1b57..08dcba2cae 100644 --- a/src/common/settings_enums.h +++ b/src/common/settings_enums.h @@ -157,6 +157,7 @@ ENUM(GpuClock, Normal, Boost, Overclock) ENUM(GpuUnswizzleSize, VerySmall, Small, Normal, Large, VeryLarge) ENUM(GpuUnswizzle, VeryLow, Low, Normal, Medium, High, Off) ENUM(GpuUnswizzleChunk, VeryLow, Low, Normal, Medium, High, Off) +ENUM(AsyncUnswizzleMode, Off, Gpu, Cpu) ENUM(TemperatureUnits, Celsius, Fahrenheit) ENUM(ExtendedDynamicState, Disabled, EDS1, EDS2, EDS3); ENUM(GpuLogLevel, Off, Errors, Standard, Verbose, All) diff --git a/src/qt_common/config/shared_translation.cpp b/src/qt_common/config/shared_translation.cpp index f086e72a7d..acf16eafaa 100644 --- a/src/qt_common/config/shared_translation.cpp +++ b/src/qt_common/config/shared_translation.cpp @@ -171,8 +171,8 @@ std::unique_ptr InitializeTranslations(QObject* parent) { "In most cases, GPU decoding provides the best performance.")); INSERT(Settings, accelerate_unswizzle, tr("Texture Unsiwzzle Method:"), tr("This option controls how generic textures should be unswizzled.\n" - "CPU: Use the CPU for unswizzling.\n" - "GPU: Use the GPU's compute shaders to unswizzling generic textures (recommended).")); + "CPU: Use the CPU for unswizzling (recommended).\n" + "GPU: Use the GPU's compute shaders to unswizzling generic textures.")); INSERT(Settings, accelerate_astc, tr("ASTC Decoding Method:"), tr("This option controls how ASTC textures should be decoded.\n" "CPU: Use the CPU for decoding.\n" @@ -234,9 +234,9 @@ std::unique_ptr InitializeTranslations(QObject* parent) { INSERT(Settings, gpu_clock, tr("GPU Clocks"), tr("Makes the game believe GPU work finishes faster than it does, so it stops lowering " "resolution and render distance to fit the Switch's clocks.")); - INSERT(Settings, gpu_unswizzle_enabled, tr("Chunked GPU Unswizzle"), + /*INSERT(Settings, gpu_unswizzle_enabled, tr("Chunked GPU Unswizzle"), tr("Accelerates BCn 3D texture decoding using GPU compute.\n" - "Disable if experiencing crashes or graphical glitches.")); + "Disable if experiencing crashes or graphical glitches."));*/ INSERT(Settings, gpu_unswizzle_texture_size, tr("GPU Unswizzle Max Texture Size"), tr("Sets the maximum size (MiB) for GPU-based texture unswizzling.\n" "While the GPU is faster for medium and large textures, the CPU may be more " @@ -250,6 +250,12 @@ std::unique_ptr InitializeTranslations(QObject* parent) { tr("Determines the number of depth slices processed in a single dispatch.\n" "Increasing this can improve throughput on high-end GPUs but may cause TDR or driver " "timeouts on weaker hardware.")); + INSERT(Settings, async_unswizzle_mode, tr("Async Unswizzle Mode"), + tr("This option controls how chunked texture uploads are unswizzled.\n" + "Off: Disables asynchronous texture unswizzling.\n" + "CPU: Uses the CPU for asynchronous unswizzling.\n" + "GPU: Accelerates BCn 3D texture decoding using GPU compute (recommended).\n" + "Change to CPU or Off if VRAM is limited.")); INSERT(Settings, use_vulkan_driver_pipeline_cache, tr("Use Vulkan pipeline cache"), tr("Enables GPU vendor-specific pipeline cache.\nThis option can improve shader loading " @@ -688,6 +694,12 @@ std::unique_ptr ComboboxEnumeration(QObject* parent) { PAIR(GpuUnswizzleChunk, High, tr("High (512)")), PAIR(GpuUnswizzleChunk, Off, tr("Off")), }}); + translations->insert({Settings::EnumMetadata::Index(), + { + PAIR(AsyncUnswizzleMode, Off, tr("Off")), + PAIR(AsyncUnswizzleMode, Cpu, tr("CPU")), + PAIR(AsyncUnswizzleMode, Gpu, tr("GPU")), + }}); translations->insert({Settings::EnumMetadata::Index(), { diff --git a/src/video_core/CMakeLists.txt b/src/video_core/CMakeLists.txt index 6720195c0e..d555179969 100644 --- a/src/video_core/CMakeLists.txt +++ b/src/video_core/CMakeLists.txt @@ -393,6 +393,19 @@ endif() if (ARCHITECTURE_x86_64) target_link_libraries(video_core PUBLIC xbyak::xbyak) + + target_sources(video_core PRIVATE + textures/decoders_avx2.cpp + textures/decoders_avx2.h + textures/decoders_sse2.cpp + textures/decoders_sse2.h + ) + if (MSVC AND NOT CMAKE_CXX_COMPILER_ID STREQUAL "Clang") + set_source_files_properties(textures/decoders_avx2.cpp PROPERTIES COMPILE_OPTIONS "/arch:AVX2") + else() + set_source_files_properties(textures/decoders_sse2.cpp PROPERTIES COMPILE_OPTIONS "-msse2") + set_source_files_properties(textures/decoders_avx2.cpp PROPERTIES COMPILE_OPTIONS "-mavx2") + endif() endif() if (ARCHITECTURE_x86_64 OR ARCHITECTURE_arm64) diff --git a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp index 17fcbb8c46..2e29c3b936 100644 --- a/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp +++ b/src/video_core/host_shaders/block_linear_unswizzle_3d_bcn.comp @@ -33,7 +33,7 @@ layout(push_constant) uniform PushConstants { uint block_depth; // Offset 48 uint block_depth_mask; // Offset 52 - int _pad; // Offset 56 + uint src_buffer_byte_bias; // Offset 56 ivec3 destination; // Offset 60 } pc; @@ -125,6 +125,7 @@ void main() { offset += (block_y & pc.block_height_mask) << GOB_SIZE_SHIFT; offset += (pos.x >> GOB_SIZE_X_SHIFT) << pc.x_shift; offset += swizzle; + offset -= pc.src_buffer_byte_bias; uvec4 texel = ReadTexel(offset); diff --git a/src/video_core/renderer_opengl/gl_texture_cache.cpp b/src/video_core/renderer_opengl/gl_texture_cache.cpp index 8ce267c16a..50699f8672 100644 --- a/src/video_core/renderer_opengl/gl_texture_cache.cpp +++ b/src/video_core/renderer_opengl/gl_texture_cache.cpp @@ -654,9 +654,11 @@ void TextureCacheRuntime::BlitFramebuffer(Framebuffer* dst, Framebuffer* src, is_linear ? GL_LINEAR : GL_NEAREST); } -void TextureCacheRuntime::AccelerateImageUpload(Image &image, const StagingBufferMap &map, +void TextureCacheRuntime::AccelerateImageUpload(Image& image, const StagingBufferMap& map, std::span swizzles, - u32 z_src_start, u32 z_image_start) { + u32 z_src_start, u32 z_image_start, u32 z_count, + [[maybe_unused]] std::span slice_has_data, + [[maybe_unused]] bool image_already_uploaded) { switch (image.info.type) { case ImageType::e2D: if (IsPixelFormatASTC(image.info.format)) { diff --git a/src/video_core/renderer_opengl/gl_texture_cache.h b/src/video_core/renderer_opengl/gl_texture_cache.h index 21a23381f4..1a3afd8956 100644 --- a/src/video_core/renderer_opengl/gl_texture_cache.h +++ b/src/video_core/renderer_opengl/gl_texture_cache.h @@ -126,9 +126,11 @@ public: const Region2D& src_region, Tegra::Engines::Fermi2D::Filter filter, Tegra::Engines::Fermi2D::Operation operation); - void AccelerateImageUpload(Image &, const StagingBufferMap &, - std::span, - u32 z_src_start, u32 z_image_start); + void AccelerateImageUpload(Image& image, const StagingBufferMap& map, + std::span swizzles, + u32 z_src_start, u32 z_image_start, u32 z_count, + std::span slice_has_data = {}, + bool image_already_uploaded = false); void InsertUploadMemoryBarrier(); diff --git a/src/video_core/renderer_vulkan/vk_compute_pass.cpp b/src/video_core/renderer_vulkan/vk_compute_pass.cpp index 345f83e53c..dded1301b2 100644 --- a/src/video_core/renderer_vulkan/vk_compute_pass.cpp +++ b/src/video_core/renderer_vulkan/vk_compute_pass.cpp @@ -1021,7 +1021,7 @@ struct alignas(16) BlockLinearUnswizzle3DPushConstants { u32 block_depth; // Offset 48 u32 block_depth_mask; // Offset 52 - s32 _pad; // Offset 56 + u32 src_buffer_byte_bias; // Offset 56 s32 destination[3]; // Offset 60 s32 _pad_end; // Offset 72 @@ -1092,18 +1092,14 @@ void BlockLinearUnswizzle3DPass::Unswizzle( std::span swizzles, u32 z_src_start, u32 z_image_start, u32 z_count, std::span slice_has_data, - std::span slice_bounds, bool image_already_uploaded) { using namespace VideoCommon::Accelerated; - const u32 MAX_BATCH_SLICES = (std::clamp)(z_count, 1u, image.info.size.depth); - static constexpr u32 MAX_WINDOW_SLICES = 4; - static constexpr float AREA_GROWTH_LIMIT = 3.0f; + const u32 MAX_BATCH_SLICES = (std::min)(z_count, image.info.size.depth); // Removing the if (!image.has_compute_unswizzle_buffer) check here is not ideal but MAX_BATCH_SLICES can changed mid-way through and I don't want to cause device loss or corruption - // This may cause issues so needs thorough testing, if works fine this is a win to reduce vram usage - image.AllocateComputeUnswizzleBuffer((std::min)(MAX_BATCH_SLICES, MAX_WINDOW_SLICES)); + image.AllocateComputeUnswizzleBuffer(MAX_BATCH_SLICES); ASSERT(swizzles.size() == 1); const auto& sw = swizzles[0]; @@ -1118,7 +1114,10 @@ void BlockLinearUnswizzle3DPass::Unswizzle( scheduler.RequestOutsideRenderPassOperationContext(); - scheduler.Record([dst_image = image.Handle(), aspect = image.AspectMask(), initial_prior_layout, level = sw.level](vk::CommandBuffer cmdbuf) { + const u32 src_buffer_byte_bias = (z_src_start < image.slice_offsets.size()) + ? image.slice_offsets[z_src_start] : 0; + + scheduler.Record([dst_image = image.Handle(), aspect = image.AspectMask(), initial_prior_layout](vk::CommandBuffer cmdbuf) { if (dst_image == VK_NULL_HANDLE) { return; } @@ -1138,126 +1137,39 @@ void BlockLinearUnswizzle3DPass::Unswizzle( .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, .image = dst_image, - .subresourceRange = {aspect, static_cast(level), 1, 0, 1}, + .subresourceRange = {aspect, 0, 1, 0, 1}, }; cmdbuf.PipelineBarrier(src_stage, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, pre_barrier); }); - struct Run { bool has_data; u32 start; u32 len; }; - boost::container::small_vector runs; - { - u32 pos = 0; - while (pos < z_count) { - const u32 src = z_src_start + pos; - const bool val = slice_has_data.empty() || - (src < static_cast(slice_has_data.size()) && slice_has_data[src] != 0); - u32 len = 1; - while (pos + len < z_count) { - const u32 next_src = z_src_start + pos + len; - const bool next_val = slice_has_data.empty() || - (next_src < static_cast(slice_has_data.size()) && slice_has_data[next_src] != 0); - if (next_val != val) break; - ++len; - } - runs.push_back({val, pos, len}); - pos += len; - } - } - - for (const Run& r : runs) { - u32 sub_offset = 0; - while (sub_offset < r.len) { - const u32 sub_len = (std::min)(r.len - sub_offset, MAX_BATCH_SLICES); - const u32 z_src = z_src_start + r.start + sub_offset; - const u32 z_dst = z_image_start + r.start + sub_offset; - - if (!r.has_data) { - // Uncomment if junk data appears - //UnswizzleZeroChunk(image, z_dst, sub_len); - sub_offset += sub_len; - continue; - } - - // Uncomment if junk data appears - /*UnswizzleZeroChunk(image, z_dst, sub_len); - - scheduler.Record([dst_image = image.Handle(), aspect = image.AspectMask()](vk::CommandBuffer cmdbuf) { - if (dst_image == VK_NULL_HANDLE) return; - const VkImageMemoryBarrier barrier{ - .sType = VK_STRUCTURE_TYPE_IMAGE_MEMORY_BARRIER, - .pNext = nullptr, - .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, - .dstAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, - .oldLayout = VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, - .newLayout = VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, - .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, - .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, - .image = dst_image, - .subresourceRange = {aspect, 0, 1, 0, 1}, - }; - cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_TRANSFER_BIT, - VK_PIPELINE_STAGE_TRANSFER_BIT, 0, barrier); - });*/ - - u32 win_offset = 0; - while (win_offset < sub_len) { - const u32 wz_start = z_src + win_offset; - const u32 wz_dst = z_dst + win_offset; - - u32 win_len = 0; - u32 ux0 = blocks_x, uy0 = blocks_y, ux1 = 0, uy1 = 0; - u64 slice_area_sum = 0; - bool any_in_window = false; - - while (win_offset + win_len < sub_len && win_len < MAX_WINDOW_SLICES) { - const u32 z = wz_start + win_len; - const bool have_box = !slice_bounds.empty() && - z < static_cast(slice_bounds.size()); - const auto b = have_box ? slice_bounds[z] : SliceBBox{}; - const bool slice_populated = b.x1 > b.x0 && b.y1 > b.y0; - - if (slice_populated) { - const u32 cx0 = any_in_window ? (std::min)(ux0, b.x0) : b.x0; - const u32 cy0 = any_in_window ? (std::min)(uy0, b.y0) : b.y0; - const u32 cx1 = any_in_window ? (std::max)(ux1, b.x1) : b.x1; - const u32 cy1 = any_in_window ? (std::max)(uy1, b.y1) : b.y1; - const u64 candidate_area = - static_cast(cx1 - cx0) * (cy1 - cy0); - const u64 candidate_slice_sum = slice_area_sum + - static_cast(b.x1 - b.x0) * (b.y1 - b.y0); - - if (any_in_window && static_cast(candidate_area) > (AREA_GROWTH_LIMIT * static_cast(candidate_slice_sum))) { - break; - } - - ux0 = cx0; uy0 = cy0; ux1 = cx1; uy1 = cy1; - slice_area_sum = candidate_slice_sum; - any_in_window = true; - } - ++win_len; - } - if (win_len == 0) { - win_len = 1; + for (u32 z_offset = 0; z_offset < z_count; z_offset += MAX_BATCH_SLICES) { + const u32 current_chunk_slices = (std::min)(MAX_BATCH_SLICES, z_count - z_offset); + const u32 current_z_src = z_src_start + z_offset; + const u32 current_z_dst = z_image_start + z_offset; + + bool chunk_has_data = slice_has_data.empty(); + if (!chunk_has_data) { + const u32 z_src_end = current_z_src + current_chunk_slices; + for (u32 z = current_z_src; z < z_src_end; ++z) { + if (z < static_cast(slice_has_data.size()) && slice_has_data[z] != 0) { + chunk_has_data = true; + break; } - - const u32 ox0 = any_in_window ? ux0 : 0; - const u32 oy0 = any_in_window ? uy0 : 0; - const u32 ox1 = any_in_window ? ux1 : blocks_x; - const u32 oy1 = any_in_window ? uy1 : blocks_y; - - UnswizzleChunk(image, swizzled, sw, params, - ox0, oy0, ox1 - ox0, oy1 - oy0, - wz_start, wz_dst, win_len); - - win_offset += win_len; } - - sub_offset += sub_len; } + + if (chunk_has_data) { + UnswizzleChunk(image, swizzled, sw, params, blocks_x, blocks_y, + current_z_src, current_z_dst, current_chunk_slices, + src_buffer_byte_bias); + }/* else { // If junk data appears uncomment this + UnswizzleZeroChunk(image, blocks_x, blocks_y, bytes_per_block, + current_z_dst, current_chunk_slices); + }*/ } - scheduler.Record([dst_image = image.Handle(), aspect = image.AspectMask(), level = sw.level](vk::CommandBuffer cmdbuf) { + scheduler.Record([dst_image = image.Handle(), aspect = image.AspectMask()](vk::CommandBuffer cmdbuf) { if (dst_image == VK_NULL_HANDLE) return; const VkImageMemoryBarrier post_barrier{ @@ -1270,7 +1182,7 @@ void BlockLinearUnswizzle3DPass::Unswizzle( .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, .image = dst_image, - .subresourceRange = {aspect, static_cast(level), 1, 0, 1}, + .subresourceRange = {aspect, 0, 1, 0, 1}, }; cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_TRANSFER_BIT, @@ -1284,16 +1196,14 @@ void BlockLinearUnswizzle3DPass::UnswizzleChunk( const StagingBufferRef& swizzled, const VideoCommon::SwizzleParameters& sw, const BlockLinearSwizzle3DParams& params, - u32 origin_x, u32 origin_y, - u32 extent_x, u32 extent_y, - u32 z_src, u32 z_dst, u32 z_count) + u32 blocks_x, u32 blocks_y, + u32 z_src, u32 z_dst, u32 z_count, + u32 src_buffer_byte_bias) { - image.compute_unswizzle_buffer_is_zero = false; - BlockLinearUnswizzle3DPushConstants pc{}; - pc.origin[0] = params.origin[0] + origin_x * 4u; - pc.origin[1] = params.origin[1] + origin_y * 4u; + pc.origin[0] = params.origin[0]; + pc.origin[1] = params.origin[1]; pc.origin[2] = z_src; // Current chunk's Z start pc.destination[0] = params.destination[0]; @@ -1308,9 +1218,10 @@ void BlockLinearUnswizzle3DPass::UnswizzleChunk( pc.block_height_mask = params.block_height_mask; pc.block_depth = params.block_depth; pc.block_depth_mask = params.block_depth_mask; + pc.src_buffer_byte_bias = src_buffer_byte_bias; - pc.blocks_dim[0] = extent_x; - pc.blocks_dim[1] = extent_y; + pc.blocks_dim[0] = blocks_x; + pc.blocks_dim[1] = blocks_y; pc.blocks_dim[2] = z_count; // Only process the count compute_pass_descriptor_queue.Acquire(scheduler, 3); @@ -1323,32 +1234,25 @@ void BlockLinearUnswizzle3DPass::UnswizzleChunk( const void* descriptor_data = compute_pass_descriptor_queue.UpdateData(); const VkDescriptorSet set = descriptor_allocator.Commit(); - const u32 gx = Common::DivCeil(extent_x, 8u); - const u32 gy = Common::DivCeil(extent_y, 8u); + const u32 gx = Common::DivCeil(blocks_x, 8u); + const u32 gy = Common::DivCeil(blocks_y, 8u); const u32 gz = Common::DivCeil(z_count, 4u); const u32 bytes_per_block = 1u << pc.bytes_per_block_log2; const VkDeviceSize output_slice_size = - static_cast(extent_x) * extent_y * bytes_per_block; + static_cast(blocks_x) * blocks_y * bytes_per_block; const VkDeviceSize barrier_size = output_slice_size * z_count; const VkBuffer out_buffer = *image.compute_unswizzle_buffer; const VkImage dst_image = image.Handle(); const VkImageAspectFlags aspect = image.AspectMask(); - const u32 level = sw.level; - const s32 dst_x = static_cast(origin_x * 4u); - const s32 dst_y = static_cast(origin_y * 4u); - - const VideoCommon::Extent3D level_size = VideoCommon::MipSize(image.info.size, level); - const u32 level_width = (std::max)(level_size.width, 1u); - const u32 level_height = (std::max)(level_size.height, 1u); - const u32 copy_width = (std::min)(extent_x * 4u, level_width - (std::min)(level_width, static_cast(dst_x))); - const u32 copy_height = (std::min)(extent_y * 4u, level_height - (std::min)(level_height, static_cast(dst_y))); + const u32 image_width = image.info.size.width; + const u32 image_height = image.info.size.height; scheduler.Record([this, set, descriptor_data, pc, gx, gy, gz, z_dst, z_count, barrier_size, - out_buffer, dst_image, aspect, level, - dst_x, dst_y, copy_width, copy_height](vk::CommandBuffer cmdbuf) { + out_buffer, dst_image, aspect, + image_width, image_height](vk::CommandBuffer cmdbuf) { if (dst_image == VK_NULL_HANDLE || out_buffer == VK_NULL_HANDLE) { return; @@ -1386,6 +1290,7 @@ void BlockLinearUnswizzle3DPass::UnswizzleChunk( .offset = 0, .size = barrier_size, }; + cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, VK_PIPELINE_STAGE_TRANSFER_BIT, 0, buffer_barrier); // Copy chunk to correct Z position in image @@ -1393,9 +1298,9 @@ void BlockLinearUnswizzle3DPass::UnswizzleChunk( .bufferOffset = 0, // Read from start of staging buffer .bufferRowLength = 0, .bufferImageHeight = 0, - .imageSubresource = {aspect, level, 0, 1}, - .imageOffset = {dst_x, dst_y, static_cast(z_dst)}, - .imageExtent = {copy_width, copy_height, z_count}, + .imageSubresource = {aspect, 0, 0, 1}, + .imageOffset = {0, 0, static_cast(z_dst)}, // Write to correct Z + .imageExtent = {image_width, image_height, z_count}, }; cmdbuf.CopyBufferToImage(out_buffer, dst_image, VK_IMAGE_LAYOUT_TRANSFER_DST_OPTIMAL, copy); @@ -1502,6 +1407,8 @@ void MSAACopyPass::CopyImage(Image& dst_image, Image& src_image, // So enjoy this mess void BlockLinearUnswizzle3DPass::UnswizzleZeroChunk( Image& image, + u32 blocks_x, u32 blocks_y, + u32 bytes_per_block, u32 z_dst, u32 z_count) { ASSERT(image.has_compute_unswizzle_buffer); @@ -1511,42 +1418,40 @@ void BlockLinearUnswizzle3DPass::UnswizzleZeroChunk( const VkImageAspectFlags aspect = image.AspectMask(); const u32 image_width = image.info.size.width; const u32 image_height = image.info.size.height; - const bool needs_fill = !image.compute_unswizzle_buffer_is_zero; - const VkDeviceSize buffer_capacity = image.compute_unswizzle_buffer_size; - if (needs_fill) { - image.compute_unswizzle_buffer_is_zero = true; - } + // Size of one unswizzled z-slice in the output buffer (bytes). + // bytes_per_block was removed here at one point but caused graphics corruption which makes sense as this is processing DXT1-7 textures and without it I'll be initilizing a buffer that is far far smaller than the actual texture + // I can look more into this later if at some point I want this to work with non-DXT textures + const VkDeviceSize output_slice_bytes = + static_cast(blocks_x) * blocks_y * bytes_per_block; + const VkDeviceSize fill_size = output_slice_bytes * z_count; scheduler.Record([out_buffer, dst_image, aspect, z_dst, z_count, - buffer_capacity, needs_fill, - image_width, image_height](vk::CommandBuffer cmdbuf) { + fill_size, image_width, image_height](vk::CommandBuffer cmdbuf) { if (dst_image == VK_NULL_HANDLE || out_buffer == VK_NULL_HANDLE) { return; } - if (needs_fill) { - // Zero the output buffer region that CopyBufferToImage will read. - cmdbuf.FillBuffer(out_buffer, 0, buffer_capacity, 0u); + // Zero the output buffer region that CopyBufferToImage will read. + cmdbuf.FillBuffer(out_buffer, 0, fill_size, 0u); - const VkBufferMemoryBarrier buffer_barrier{ - .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, - .pNext = nullptr, - .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, - .dstAccessMask = VK_ACCESS_TRANSFER_READ_BIT, - .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, - .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, - .buffer = out_buffer, - .offset = 0, - .size = buffer_capacity, - }; + const VkBufferMemoryBarrier buffer_barrier{ + .sType = VK_STRUCTURE_TYPE_BUFFER_MEMORY_BARRIER, + .pNext = nullptr, + .srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT, + .dstAccessMask = VK_ACCESS_TRANSFER_READ_BIT, + .srcQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .dstQueueFamilyIndex = VK_QUEUE_FAMILY_IGNORED, + .buffer = out_buffer, + .offset = 0, + .size = fill_size, + }; - cmdbuf.PipelineBarrier( - VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, - VK_PIPELINE_STAGE_TRANSFER_BIT, - 0, buffer_barrier); - } + cmdbuf.PipelineBarrier( + VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, + VK_PIPELINE_STAGE_TRANSFER_BIT, + 0, buffer_barrier); // Copy the zeroed buffer region into the correct Z position of the image. const VkBufferImageCopy copy{ diff --git a/src/video_core/renderer_vulkan/vk_compute_pass.h b/src/video_core/renderer_vulkan/vk_compute_pass.h index 5a2253da21..e7f344b467 100644 --- a/src/video_core/renderer_vulkan/vk_compute_pass.h +++ b/src/video_core/renderer_vulkan/vk_compute_pass.h @@ -150,7 +150,6 @@ public: std::span swizzles, u32 z_src_start, u32 z_image_start, u32 z_count, std::span slice_has_data, - std::span slice_bounds, bool image_already_uploaded); void UnswizzleChunk( @@ -158,12 +157,14 @@ public: const StagingBufferRef &swizzled, const VideoCommon::SwizzleParameters &sw, const BlockLinearSwizzle3DParams ¶ms, - u32 origin_x, u32 origin_y, - u32 extent_x, u32 extent_y, - u32 z_src, u32 z_dst, u32 z_count); + u32 blocks_x, u32 blocks_y, + u32 z_src, u32 z_dst, u32 z_count, + u32 src_buffer_byte_bias); void UnswizzleZeroChunk( Image &image, + u32 blocks_x, u32 blocks_y, + u32 bytes_per_block, u32 z_dst, u32 z_count); private: diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.cpp b/src/video_core/renderer_vulkan/vk_texture_cache.cpp index 58b20e5dec..fd9c9cc45a 100644 --- a/src/video_core/renderer_vulkan/vk_texture_cache.cpp +++ b/src/video_core/renderer_vulkan/vk_texture_cache.cpp @@ -949,13 +949,14 @@ void BlitScale(Scheduler& scheduler, VkImage src_image, VkImage dst_image, const } } -[[nodiscard]] boost::container::small_vector BuildViewFormats( - const ImageInfo& info, std::span base_view_formats) { +[[nodiscard]] boost::container::small_vector BuildViewFormats(const ImageInfo& info, std::span base_view_formats) { boost::container::small_vector formats(base_view_formats.begin(), base_view_formats.end()); - if (Settings::values.gpu_unswizzle_enabled.GetValue()) { - if (const auto block_view_format = BlockTexelViewFormat(info.format)) { - formats.push_back(*block_view_format); + if (Settings::values.accelerate_unswizzle.GetValue() == Settings::TexUnswizzleMode::Gpu) { + if (!info.is_sparse) { + if (const auto block_view_format = BlockTexelViewFormat(info.format)) { + formats.push_back(*block_view_format); + } } } return formats; @@ -999,11 +1000,7 @@ TextureCacheRuntime::TextureCacheRuntime(const Device& device_, Scheduler& sched } } - if (Settings::values.gpu_unswizzle_enabled.GetValue()) { - bl2d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, - staging_buffer_pool, compute_pass_descriptor_queue); - bl3d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, - staging_buffer_pool, compute_pass_descriptor_queue); + if (Settings::values.accelerate_unswizzle.GetValue() == Settings::TexUnswizzleMode::Gpu) { generic_2d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, staging_buffer_pool, compute_pass_descriptor_queue); generic_3d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, @@ -1011,6 +1008,13 @@ TextureCacheRuntime::TextureCacheRuntime(const Device& device_, Scheduler& sched generic_linear_unswizzle_pass.emplace(device, scheduler, descriptor_pool, staging_buffer_pool, compute_pass_descriptor_queue); } + + if (Settings::values.async_unswizzle_mode.GetValue() == Settings::AsyncUnswizzleMode::Gpu) { + bl2d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, + staging_buffer_pool, compute_pass_descriptor_queue); + bl3d_unswizzle_pass.emplace(device, scheduler, descriptor_pool, + staging_buffer_pool, compute_pass_descriptor_queue); + } } bool TextureCacheRuntime::IsUnswizzleStorageFormatSupported(PixelFormat format) const { @@ -2989,17 +2993,24 @@ VkRenderPass Framebuffer::RenderPassVariant(u32 color_clear_mask, bool depth_ste void TextureCacheRuntime::AccelerateImageUpload( Image& image, const StagingBufferRef& map, std::span swizzles, - u32 z_src_start, u32 z_image_start) { + u32 z_src_start, u32 z_image_start, u32 z_count, + std::span slice_has_data, + bool image_already_uploaded) { if (IsPixelFormatASTC(image.info.format)) { return astc_decoder_pass->Assemble(image, map, swizzles); } - if (!Settings::values.gpu_unswizzle_enabled.GetValue() || (!generic_2d_unswizzle_pass && !generic_3d_unswizzle_pass)) { + /*if (!Settings::values.gpu_unswizzle_enabled.GetValue() || (!generic_2d_unswizzle_pass && !generic_3d_unswizzle_pass)) { ASSERT(false && "GPU unswizzle is disabled for this texture"); return; - } + }*/ + if (bl3d_unswizzle_pass && IsPixelFormatBCn(image.info.format) && image.info.type == ImageType::e3D && image.info.resources.levels == 1 && image.info.resources.layers == 1) { + return bl3d_unswizzle_pass->Unswizzle(image, map, swizzles, + z_src_start, z_image_start, z_count, + slice_has_data, image_already_uploaded); + } if (image.info.type == ImageType::e2D && generic_2d_unswizzle_pass) { return generic_2d_unswizzle_pass->Unswizzle(image, map, swizzles); } diff --git a/src/video_core/renderer_vulkan/vk_texture_cache.h b/src/video_core/renderer_vulkan/vk_texture_cache.h index 475ef83c90..a6112c5130 100644 --- a/src/video_core/renderer_vulkan/vk_texture_cache.h +++ b/src/video_core/renderer_vulkan/vk_texture_cache.h @@ -101,9 +101,11 @@ public: Settings::TexUnswizzleMode::Gpu; } - void AccelerateImageUpload(Image &, const StagingBufferRef &, + void AccelerateImageUpload(Image&, const StagingBufferRef&, std::span, - u32 z_src_start, u32 z_image_start); + u32 z_src_start, u32 z_image_start, u32 z_count, + std::span slice_has_data = {}, + bool image_already_uploaded = false); void InsertUploadMemoryBarrier() {} diff --git a/src/video_core/texture_cache/accelerated_swizzle.h b/src/video_core/texture_cache/accelerated_swizzle.h index 496a1791d7..7de17681b2 100644 --- a/src/video_core/texture_cache/accelerated_swizzle.h +++ b/src/video_core/texture_cache/accelerated_swizzle.h @@ -14,9 +14,6 @@ namespace VideoCommon::Accelerated { -constexpr u32 GOB_SIZE_Y = 8; -constexpr u32 GOB_SIZE = 512; - struct BlockLinearSwizzle2DParams { alignas(16) std::array origin; alignas(16) std::array destination; @@ -41,85 +38,10 @@ struct BlockLinearSwizzle3DParams { u32 block_depth_mask; }; -struct SliceBBox { - u32 x0 = 0, y0 = 0, x1 = 0, y1 = 0; -}; - [[nodiscard]] BlockLinearSwizzle2DParams MakeBlockLinearSwizzle2DParams( const SwizzleParameters& swizzle, const ImageInfo& info); [[nodiscard]] BlockLinearSwizzle3DParams MakeBlockLinearSwizzle3DParams( const SwizzleParameters& swizzle, const ImageInfo& info); -template -void ForEachZInGroupOverlap(u64 local_start, u64 local_end, - u32 block_size, u32 x_shift, u32 block_height, - u32 block_height_mask, u32 block_depth, u32 block_depth_mask, - u32 bytes_per_block, u32 blocks_x, u32 blocks_y, - Callback&& on_result) { - if (local_start >= local_end) return; - - const u32 z_pitch = GOB_SIZE << block_height; - const u32 column_pitch = 1u << x_shift; - const u32 blocks_per_column = 64u / bytes_per_block; - const u32 rows_per_group = 1u << block_height; - const u32 y_per_band = GOB_SIZE_Y << block_height; - const u32 zz_count = 1u << block_depth; - - const u64 band_start = local_start / block_size; - const u64 band_end = (local_end - 1) / block_size; - - if (band_start != band_end) { - SliceBBox box; - box.x0 = 0; - box.x1 = blocks_x; - box.y0 = static_cast(band_start) * y_per_band; - box.y1 = (std::min)(static_cast(band_end + 1) * y_per_band, blocks_y); - for (u32 zz = 0; zz < zz_count; ++zz) on_result(zz, box); - return; - } - - const u64 wb_start = local_start - band_start * block_size; - const u64 wb_end = (local_end - 1) - band_start * block_size; - const u32 col_start = static_cast(wb_start / column_pitch); - const u32 col_end = static_cast(wb_end / column_pitch); - - if (col_start != col_end) { - SliceBBox box; - box.x0 = col_start * blocks_per_column; - box.x1 = (std::min)((col_end + 1) * blocks_per_column, blocks_x); - box.y0 = static_cast(band_start) * y_per_band; - box.y1 = (std::min)(box.y0 + y_per_band, blocks_y); - for (u32 zz = 0; zz < zz_count; ++zz) on_result(zz, box); - return; - } - - const u64 wc_start = wb_start - static_cast(col_start) * column_pitch; - const u64 wc_end = wb_end - static_cast(col_start) * column_pitch; - const u32 zz_start = static_cast(wc_start / z_pitch); - const u32 zz_end = static_cast(wc_end / z_pitch); - - if (zz_start != zz_end) { - SliceBBox box; - box.x0 = col_start * blocks_per_column; - box.x1 = (std::min)((col_start + 1) * blocks_per_column, blocks_x); - box.y0 = static_cast(band_start) * y_per_band; - box.y1 = (std::min)(box.y0 + y_per_band, blocks_y); - for (u32 zz = zz_start; zz <= zz_end; ++zz) on_result(zz, box); - return; - } - - const u64 wz_start = wc_start - static_cast(zz_start) * z_pitch; - const u64 wz_end = wc_end - static_cast(zz_start) * z_pitch; - const u32 row_start = static_cast(wz_start) / GOB_SIZE; - const u32 row_end = static_cast(wz_end) / GOB_SIZE; - - SliceBBox box; - box.x0 = col_start * blocks_per_column; - box.x1 = (std::min)((col_start + 1) * blocks_per_column, blocks_x); - box.y0 = (static_cast(band_start) * rows_per_group + row_start) * GOB_SIZE_Y; - box.y1 = (std::min)((static_cast(band_start) * rows_per_group + row_end + 1) * GOB_SIZE_Y, blocks_y); - on_result(zz_start, box); -} - } // namespace VideoCommon::Accelerated diff --git a/src/video_core/texture_cache/texture_cache.h b/src/video_core/texture_cache/texture_cache.h index 62a1edb990..69e2a68670 100644 --- a/src/video_core/texture_cache/texture_cache.h +++ b/src/video_core/texture_cache/texture_cache.h @@ -31,6 +31,7 @@ namespace VideoCommon { using Tegra::Texture::TICEntry; using Tegra::Texture::TSCEntry; +using Tegra::Texture::UnswizzleTexture; using VideoCore::Surface::GetFormatType; using VideoCore::Surface::PixelFormat; using VideoCore::Surface::SurfaceType; @@ -78,41 +79,41 @@ TextureCache

::TextureCache(Runtime& runtime_, Tegra::MaxwellDeviceMemoryManag minimum_memory = 0; } - const bool gpu_unswizzle_enabled = Settings::values.gpu_unswizzle_enabled.GetValue(); + async_unswizzle_mode = Settings::values.async_unswizzle_mode.GetValue(); - if (gpu_unswizzle_enabled) { + if (async_unswizzle_mode != Settings::AsyncUnswizzleMode::Off) { switch (Settings::values.gpu_unswizzle_texture_size.GetValue()) { - case Settings::GpuUnswizzleSize::VerySmall: gpu_unswizzle_maxsize = 16_MiB; break; - case Settings::GpuUnswizzleSize::Small: gpu_unswizzle_maxsize = 32_MiB; break; - case Settings::GpuUnswizzleSize::Normal: gpu_unswizzle_maxsize = 128_MiB; break; - case Settings::GpuUnswizzleSize::Large: gpu_unswizzle_maxsize = 256_MiB; break; - case Settings::GpuUnswizzleSize::VeryLarge: gpu_unswizzle_maxsize = 512_MiB; break; - default: gpu_unswizzle_maxsize = 128_MiB; break; + case Settings::GpuUnswizzleSize::VerySmall: async_unswizzle_maxsize = 16_MiB; break; + case Settings::GpuUnswizzleSize::Small: async_unswizzle_maxsize = 32_MiB; break; + case Settings::GpuUnswizzleSize::Normal: async_unswizzle_maxsize = 128_MiB; break; + case Settings::GpuUnswizzleSize::Large: async_unswizzle_maxsize = 256_MiB; break; + case Settings::GpuUnswizzleSize::VeryLarge: async_unswizzle_maxsize = 512_MiB; break; + default: async_unswizzle_maxsize = 128_MiB; break; } switch (Settings::values.gpu_unswizzle_stream_size.GetValue()) { - case Settings::GpuUnswizzle::VeryLow: swizzle_chunk_size = 4_MiB; break; - case Settings::GpuUnswizzle::Low: swizzle_chunk_size = 8_MiB; break; - case Settings::GpuUnswizzle::Normal: swizzle_chunk_size = 16_MiB; break; - case Settings::GpuUnswizzle::Medium: swizzle_chunk_size = 32_MiB; break; - case Settings::GpuUnswizzle::High: swizzle_chunk_size = 64_MiB; break; - case Settings::GpuUnswizzle::Off: swizzle_chunk_size = 0; break; - default: swizzle_chunk_size = 16_MiB; + case Settings::GpuUnswizzle::VeryLow: async_unswizzle_chunk_size = 4_MiB; break; + case Settings::GpuUnswizzle::Low: async_unswizzle_chunk_size = 8_MiB; break; + case Settings::GpuUnswizzle::Normal: async_unswizzle_chunk_size = 16_MiB; break; + case Settings::GpuUnswizzle::Medium: async_unswizzle_chunk_size = 32_MiB; break; + case Settings::GpuUnswizzle::High: async_unswizzle_chunk_size = 64_MiB; break; + case Settings::GpuUnswizzle::Off: async_unswizzle_chunk_size = 0; break; + default: async_unswizzle_chunk_size = 16_MiB; } switch (Settings::values.gpu_unswizzle_chunk_size.GetValue()) { - case Settings::GpuUnswizzleChunk::VeryLow: swizzle_slices_per_batch = 32; break; - case Settings::GpuUnswizzleChunk::Low: swizzle_slices_per_batch = 64; break; - case Settings::GpuUnswizzleChunk::Normal: swizzle_slices_per_batch = 128; break; - case Settings::GpuUnswizzleChunk::Medium: swizzle_slices_per_batch = 256; break; - case Settings::GpuUnswizzleChunk::High: swizzle_slices_per_batch = 512; break; - case Settings::GpuUnswizzleChunk::Off: swizzle_slices_per_batch = 0; break; - default: swizzle_slices_per_batch = 128; + case Settings::GpuUnswizzleChunk::VeryLow: async_unswizzle_slices_per_batch = 32; break; + case Settings::GpuUnswizzleChunk::Low: async_unswizzle_slices_per_batch = 64; break; + case Settings::GpuUnswizzleChunk::Normal: async_unswizzle_slices_per_batch = 128; break; + case Settings::GpuUnswizzleChunk::Medium: async_unswizzle_slices_per_batch = 256; break; + case Settings::GpuUnswizzleChunk::High: async_unswizzle_slices_per_batch = 512; break; + case Settings::GpuUnswizzleChunk::Off: async_unswizzle_slices_per_batch = 0; break; + default: async_unswizzle_slices_per_batch = 128; } } else { - gpu_unswizzle_maxsize = 0; - swizzle_chunk_size = 0; - swizzle_slices_per_batch = 0; + async_unswizzle_maxsize = 0; + async_unswizzle_chunk_size = 0; + async_unswizzle_slices_per_batch = 0; } } @@ -1119,17 +1120,18 @@ void TextureCache

::RefreshContents(Image& image, ImageId image_id) { return; } - if (Settings::values.gpu_unswizzle_enabled.GetValue() && + if (async_unswizzle_mode != Settings::AsyncUnswizzleMode::Off && IsPixelFormatBCn(image.info.format) && image.info.type == ImageType::e3D && image.info.resources.levels == 1 && image.info.resources.layers == 1 && - MapSizeBytes(image) >= gpu_unswizzle_maxsize && - False(image.flags & ImageFlagBits::GpuModified)) { + False(image.flags & ImageFlagBits::GpuModified) && + MapSizeBytes(image) >= async_unswizzle_maxsize) { QueueAsyncUnswizzle(image, image_id); return; } + auto staging = runtime.UploadStagingBuffer(MapSizeBytes(image)); UploadImageContents(image, staging); runtime.InsertUploadMemoryBarrier(); @@ -1146,7 +1148,7 @@ void TextureCache

::UploadImageContents(Image& image, StagingBuffer& staging) const auto uploads = FullUploadSwizzles(image.info); gpu_memory->ReadBlock(gpu_addr, mapped_span.data(), mapped_span.size_bytes(), VideoCommon::CacheType::NoTextureCache); - runtime.AccelerateImageUpload(image, staging, FixSmallVectorADL(uploads), 0, 0); + runtime.AccelerateImageUpload(image, staging, FixSmallVectorADL(uploads), 0, 0, 0); return; } } @@ -1366,10 +1368,12 @@ void TextureCache

::QueueAsyncUnswizzle(Image& image, ImageId image_id) { } image.flags |= ImageFlagBits::IsDecoding; + image.flags |= ImageFlagBits::AcceleratedUpload; unswizzle_queue.push_back({ + .info = image.info, .image_id = image_id, - .info = image.info + .is_cpu = async_unswizzle_mode == Settings::AsyncUnswizzleMode::Cpu, }); } @@ -1405,72 +1409,219 @@ void TextureCache

::TickAsyncUnswizzle() { } PendingUnswizzle& task = unswizzle_queue.front(); - Image& image = slot_images[task.image_id]; + ImageId task_image_id = task.image_id; + + if (task.is_cpu) { + // The primary resource drain here is UnswizzleTexture so maybe this won't blow up your SteamDeck + // Also, scary infinite loop possibility ... makes me uncomfortable + while (true) { + PendingUnswizzle& current = unswizzle_queue.front(); + Image& image = slot_images[task_image_id]; + const bool was_job_in_flight = current.cpu_job_in_flight; + const size_t offset_before = current.last_submitted_offset; + + if (was_job_in_flight) { + TickAsyncUnswizzleCpu(current, image); + break; + } + + TickAsyncUnswizzleCpu(current, image); + + if (unswizzle_queue.empty() || unswizzle_queue.front().image_id != task_image_id) { + break; + } + + PendingUnswizzle& refreshed = unswizzle_queue.front(); + if (refreshed.cpu_job_in_flight) { + break; + } + if (refreshed.last_submitted_offset == offset_before) { + break; + } + } + } else { + Image& image = slot_images[task_image_id]; + TickAsyncUnswizzleGpu(task, image); + } +} +template +void TextureCache

::TickAsyncUnswizzleGpu(PendingUnswizzle& task, Image& image) { if (!task.initialized) { - if (!task.is_incremental) { - task.total_size = MapSizeBytes(image); - } - task.staging_buffer = runtime.UploadStagingBuffer(task.total_size, true); - - if (!task.is_incremental) { - const auto& info = image.info; - const u32 bytes_per_block = BytesPerBlock(info.format); - const u32 width_blocks = Common::DivCeil(info.size.width, 4u); - const u32 height_blocks = Common::DivCeil(info.size.height, 4u); - task.bytes_per_slice = static_cast(width_blocks * bytes_per_block) * height_blocks; - task.last_submitted_offset = 0; - - task.is_sparse = True(image.flags & ImageFlagBits::Sparse); - if (task.is_sparse) { - const auto segs = - gpu_memory->GetSubmappedRange(image.gpu_addr, image.guest_size_bytes); - task.sparse_segments.assign(segs.begin(), segs.end()); - - task.segment_scan_cursor = 0; - task.slice_has_data.assign(image.info.size.depth, 0u); - - if (image.info.size.depth > 1) { - const auto uploads = FullUploadSwizzles(task.info); - const auto sp = VideoCommon::Accelerated::MakeBlockLinearSwizzle3DParams( - uploads[0], task.info); - - task.swizzle_group_size = sp.slice_size; - task.slices_per_group = 1u << sp.block_depth; - - const u32 num_groups = Common::DivCeil( - static_cast(image.info.size.depth), task.slices_per_group); - task.slice_has_data.assign(num_groups, 0u); - - if (task.swizzle_group_size > 0) { - for (const auto& [seg_gpu_addr, seg_size] : task.sparse_segments) { - if (seg_gpu_addr < image.gpu_addr) continue; - const u64 seg_start = seg_gpu_addr - image.gpu_addr; - const u64 seg_end = seg_start + seg_size; - const u32 g_first = static_cast(seg_start / task.swizzle_group_size); - const u32 g_last = static_cast((seg_end - 1) / task.swizzle_group_size); - for (u32 g = g_first; g <= g_last && g < num_groups; ++g) { - task.slice_has_data[g] = 1u; - } + task.total_size = MapSizeBytes(image); + + const auto& info = image.info; + const u32 bytes_per_block = BytesPerBlock(info.format); + const u32 width_blocks = Common::DivCeil(info.size.width, 4u); + const u32 height_blocks = Common::DivCeil(info.size.height, 4u); + task.bytes_per_slice = static_cast(width_blocks * bytes_per_block) * height_blocks; + task.last_submitted_offset = 0; + + const u32 total_slices_init = info.size.depth; + + task.is_sparse = True(image.flags & ImageFlagBits::Sparse); + task.active_z_start = 0; + task.active_z_end = total_slices_init; + if (task.is_sparse) { + const auto segs = + gpu_memory->GetSubmappedRange(image.gpu_addr, image.guest_size_bytes); + task.sparse_segments.assign(segs.begin(), segs.end()); + task.segment_scan_cursor = 0; + + task.slice_has_data.assign(image.info.size.depth, 0u); + + if (image.info.size.depth > 1 && !image.slice_offsets.empty()) { + const auto uploads = FullUploadSwizzles(task.info); + const auto sp = VideoCommon::Accelerated::MakeBlockLinearSwizzle3DParams( + uploads[0], task.info); + const u64 swizzled_slice_size = sp.slice_size; + task.swizzled_slice_size = swizzled_slice_size; + task.swizzle_block_depth = sp.block_depth; + + const u32 depth = static_cast(image.info.size.depth); + const u32 group_slices = 1u << task.swizzle_block_depth; + const u64 group_byte_size = task.swizzled_slice_size; + + size_t seg_idx = 0; + for (u32 g_start = 0; g_start < depth; g_start += group_slices) { + const u32 g_end = (std::min)(g_start + group_slices, depth); + const u64 group_abs_start = image.slice_offsets[g_start]; + const u64 group_abs_end = + (g_end < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[g_end]) + : group_abs_start + group_byte_size; + + while (seg_idx < task.sparse_segments.size()) { + const auto& [seg_gpu_addr, seg_size] = task.sparse_segments[seg_idx]; + const u64 seg_start = seg_gpu_addr - image.gpu_addr; + const u64 seg_end = seg_start + seg_size; + if (seg_end <= group_abs_start) { + ++seg_idx; + continue; + } + break; + } + + bool group_has_data = false; + for (size_t probe = seg_idx; probe < task.sparse_segments.size(); ++probe) { + const auto& [seg_gpu_addr, seg_size] = task.sparse_segments[probe]; + const u64 seg_start = seg_gpu_addr - image.gpu_addr; + const u64 seg_end = seg_start + seg_size; + if (seg_start >= group_abs_end) break; + if (seg_end > group_abs_start && seg_start < group_abs_end) { + group_has_data = true; + break; } } + + if (group_has_data) { + for (u32 z = g_start; z < g_end; ++z) { + task.slice_has_data[z] = 1u; + } + } + } + } else { + std::fill(task.slice_has_data.begin(), task.slice_has_data.end(), 1u); + } + + if (async_unswizzle_slices_per_batch == 0) { + u32 first_populated = total_slices_init; + u32 last_populated_excl = 0; + for (u32 z = 0; z < total_slices_init; ++z) { + if (task.slice_has_data[z]) { + first_populated = (std::min)(first_populated, z); + last_populated_excl = z + 1; + } + } + if (first_populated < last_populated_excl) { + const u32 group_slices_align = (std::max)(1u << task.swizzle_block_depth, 1u); + task.active_z_start = (first_populated / group_slices_align) * group_slices_align; + task.active_z_end = (std::min)( + Common::AlignUp(last_populated_excl, group_slices_align), total_slices_init); } else { - std::fill(task.slice_has_data.begin(), task.slice_has_data.end(), 1u); + task.active_z_start = 0; + task.active_z_end = 0; } } } + size_t max_batch_size = task.total_size; + { + const u32 batch_slices = (async_unswizzle_slices_per_batch == 0) + ? (task.active_z_end - task.active_z_start) + : async_unswizzle_slices_per_batch; + + size_t calculated_max = 0; + for (u32 z = task.active_z_start; z < task.active_z_end; z += (std::max)(batch_slices, 1u)) { + const u32 z_want_init = (std::min)(z + (std::max)(batch_slices, 1u), task.active_z_end); + const u32 z_count_init = z_want_init - z; + + const size_t group_off = (z < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z]) : 0; + const size_t group_end = (z + z_count_init < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z + z_count_init]) : task.total_size; + + calculated_max = (std::max)(calculated_max, group_end - group_off); + } + if (calculated_max > 0 && calculated_max < max_batch_size) { + max_batch_size = calculated_max; + } + if (task.active_z_start >= task.active_z_end) { + max_batch_size = (std::min)(max_batch_size, task.bytes_per_slice); + } + + if (task.active_z_start > 0 && task.active_z_start < image.slice_offsets.size()) { + task.current_offset = + static_cast(image.slice_offsets[task.active_z_start]); + task.last_submitted_offset = + static_cast(task.active_z_start) * task.bytes_per_slice; + } + } + + const size_t needed = (std::max)(max_batch_size, size_t{1}); + if (needed > UnswizzleSharedStagingCap) { + task.staging_buffer = runtime.UploadStagingBuffer(needed, true); + task.owns_staging_buffer = true; + } else { + if (unswizzle_shared_staging_pending_gpu_read) { + runtime.Finish(); + unswizzle_shared_staging_pending_gpu_read = false; + } + + if (!unswizzle_shared_staging.has_value() || + unswizzle_shared_staging_capacity < needed) { + if (unswizzle_shared_staging.has_value()) { + runtime.FreeDeferredStagingBuffer(*unswizzle_shared_staging); + } + const size_t grown = (std::min)( + UnswizzleSharedStagingCap, + (std::max)(needed, unswizzle_shared_staging_capacity + + unswizzle_shared_staging_capacity / 4)); + unswizzle_shared_staging = runtime.UploadStagingBuffer(grown, true); + unswizzle_shared_staging_capacity = grown; + } + task.staging_buffer = *unswizzle_shared_staging; + task.staging_buffer.mapped_span = + task.staging_buffer.mapped_span.subspan(0, needed); + task.owns_staging_buffer = false; + } + task.initialized = true; } // Read data - if (task.current_offset < task.total_size) { - const size_t remaining = task.total_size - task.current_offset; - size_t copy_amount = (swizzle_chunk_size == 0 || task.is_incremental) + const size_t active_start_byte = (task.active_z_start < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[task.active_z_start]) : 0; + const size_t active_end_byte = (task.active_z_end < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[task.active_z_end]) : task.total_size; + + if (task.current_offset < active_end_byte) { + const size_t remaining = active_end_byte - task.current_offset; + size_t copy_amount = (async_unswizzle_chunk_size == 0) ? remaining - : (std::min)(swizzle_chunk_size, remaining); + : (std::min)(async_unswizzle_chunk_size, remaining); - if (swizzle_chunk_size > 0 && !task.is_incremental && copy_amount < remaining) { + if (async_unswizzle_chunk_size > 0 && copy_amount < remaining) { copy_amount = (copy_amount / task.bytes_per_slice) * task.bytes_per_slice; if (copy_amount == 0) copy_amount = task.bytes_per_slice; copy_amount = (std::min)(copy_amount, remaining); @@ -1478,6 +1629,11 @@ void TextureCache

::TickAsyncUnswizzle() { u8* const staging_base = task.staging_buffer.mapped_span.data(); + if (task.is_sparse && task.current_offset == active_start_byte) { + const size_t window_size = active_end_byte - active_start_byte; + std::memset(staging_base, 0, (std::min)(window_size, task.staging_buffer.mapped_span.size())); + } + const size_t base_off = task.staging_base_byte_offset; const size_t read_start = task.current_offset; const size_t read_end = read_start + copy_amount; @@ -1485,35 +1641,6 @@ void TextureCache

::TickAsyncUnswizzle() { const size_t abs_end = read_end + base_off; if (task.is_sparse) { - size_t cursor = read_start; - - const bool can_smart_skip = - task.swizzle_group_size > 0 && - task.slices_per_group > 0 && - !task.slice_has_data.empty(); - - auto fill_gap = [&](size_t gap_start_rel, size_t gap_end_rel) { - if (gap_start_rel >= gap_end_rel) return; - if (!can_smart_skip) { - std::memset(staging_base + gap_start_rel, 0, gap_end_rel - gap_start_rel); - return; - } - size_t pos = gap_start_rel; - while (pos < gap_end_rel) { - const u64 abs_pos = pos + base_off; - const u32 z_group = static_cast(abs_pos / task.swizzle_group_size); - const u32 z_group_local = z_group - - (task.incremental_z_start / task.slices_per_group); - if (z_group_local >= static_cast(task.slice_has_data.size())) break; - const size_t group_abs_end = - (static_cast(z_group) + 1) * task.swizzle_group_size; - const size_t end_rel = (std::min)(gap_end_rel, group_abs_end - base_off); - if (task.slice_has_data[z_group_local]) - std::memset(staging_base + pos, 0, end_rel - pos); - pos = end_rel; - } - }; - while (task.segment_scan_cursor < task.sparse_segments.size()) { const auto& [seg_gpu_addr, seg_size] = task.sparse_segments[task.segment_scan_cursor]; @@ -1526,77 +1653,562 @@ void TextureCache

::TickAsyncUnswizzle() { const size_t ol_abs_start = (std::max)(seg_abs_start, abs_start); const size_t ol_abs_end = (std::min)(seg_abs_end, abs_end); - - const size_t ol_rel_start = ol_abs_start - base_off; - const size_t ol_rel_end = ol_abs_end - base_off; - - fill_gap(cursor, ol_rel_start); + const size_t ol_rel_start = (ol_abs_start - base_off) - active_start_byte; gpu_memory->ReadBlockUnsafe(image.gpu_addr + ol_abs_start, staging_base + ol_rel_start, ol_abs_end - ol_abs_start); - cursor = ol_rel_end; if (seg_abs_end > abs_end) break; ++task.segment_scan_cursor; } - fill_gap(cursor, read_end); - } else { gpu_memory->ReadBlockUnsafe(image.gpu_addr + abs_start, - staging_base + read_start, + staging_base + (read_start - active_start_byte), copy_amount); } task.current_offset += copy_amount; } - const bool is_final_batch = task.current_offset >= task.total_size; - const bool is_sparse_3d = task.is_sparse && task.swizzle_group_size > 0 && task.slices_per_group > 0; - const size_t effective_group_size = is_sparse_3d ? task.swizzle_group_size : task.bytes_per_slice; - const u32 slices_per_group_val = is_sparse_3d ? task.slices_per_group : 1u; - + const bool is_final_batch = task.current_offset >= active_end_byte; const size_t bytes_ready = task.current_offset - task.last_submitted_offset; - const u32 complete_slices = static_cast(bytes_ready / effective_group_size) * slices_per_group_val; - - const u32 total_slices = task.is_incremental - ? task.incremental_z_count - : image.info.size.depth; - const u32 batch = task.is_incremental - ? task.incremental_z_count - : (swizzle_slices_per_batch == 0 ? image.info.size.depth : swizzle_slices_per_batch); - - if (complete_slices >= slices_per_group_val || (is_final_batch && complete_slices > 0)) { - const u32 z_src = static_cast(task.last_submitted_offset / effective_group_size) * slices_per_group_val; - const u32 z_image = task.incremental_z_start + z_src; - const u32 z_count = (std::min)({complete_slices, batch, total_slices - z_src}); + const u32 complete_slices = static_cast(bytes_ready / task.bytes_per_slice); + + const std::span sparse_hint = + task.is_sparse ? std::span(task.slice_has_data) + : std::span{}; + + const u32 total_slices = image.info.size.depth; + const u32 effective_end = (async_unswizzle_slices_per_batch == 0 && task.is_sparse) + ? task.active_z_end + : total_slices; + const u32 batch = + (async_unswizzle_slices_per_batch == 0 ? effective_end : async_unswizzle_slices_per_batch); + + if (complete_slices >= batch || (is_final_batch && complete_slices > 0)) { + const u32 z_src = static_cast(task.last_submitted_offset / task.bytes_per_slice); + const u32 z_image = z_src; // + 0 for full tasks + const u32 z_count = (z_src < effective_end) + ? (std::min)({complete_slices, batch, effective_end - z_src}) + : 0; if (z_count > 0) { auto uploads = FullUploadSwizzles(task.info); - if (task.is_incremental) { - uploads[0].num_tiles.depth = task.incremental_z_count; - } else if (task.is_sparse) { - uploads[0].num_tiles.depth = z_count; + if (!task.owns_staging_buffer) { + unswizzle_shared_staging_pending_gpu_read = true; } runtime.AccelerateImageUpload(image, task.staging_buffer, FixSmallVectorADL(uploads), - z_src, z_image); - const u32 groups_dispatched = z_count / slices_per_group_val; - task.last_submitted_offset += static_cast(groups_dispatched) * effective_group_size; + z_src, z_image, z_count, + sparse_hint, + false); + task.last_submitted_offset += static_cast(z_count) * task.bytes_per_slice; } } // Check if complete - const u32 slices_submitted = static_cast(task.last_submitted_offset / effective_group_size) * slices_per_group_val; - const bool all_submitted = slices_submitted >= total_slices || - (is_final_batch && bytes_ready < effective_group_size); + const u32 slices_submitted = static_cast(task.last_submitted_offset / task.bytes_per_slice); + const bool all_submitted = slices_submitted >= effective_end || + (is_final_batch && bytes_ready < task.bytes_per_slice); if (is_final_batch && all_submitted) { - runtime.FreeDeferredStagingBuffer(task.staging_buffer); + if (task.owns_staging_buffer) { + runtime.FreeDeferredStagingBuffer(task.staging_buffer); + } runtime.ReleaseSparseUnswizzleBuffer(image); image.flags &= ~ImageFlagBits::IsDecoding; unswizzle_queue.pop_front(); } } +// Send help, I didn't think this would be that complicated to implement +// Whatever, this needlessly over complicated block of code is this way simply to save RAM +// +// Note for Liz: +// Divided everything into helper functions but don't know if this makes everything more confusing. +// Or could possibly cause a speed reduction, I'm not entirely sure what template functions look like in ASM +// So I am unsure if there will be overhead with saving the return location into the stack +template +void TextureCache

::InitializeCpuUnswizzleTask(PendingUnswizzle& task, Image& image) { + task.total_size = MapSizeBytes(image); + + const auto uploads = FullUploadSwizzles(task.info); + task.cpu_num_tiles = uploads[0].num_tiles; + task.cpu_block = uploads[0].block; + task.cpu_stride_alignment = CalculateLevelStrideAlignment(task.info, 0); + task.cpu_bytes_per_block = BytesPerBlock(task.info.format); + + size_t max_batch_size = task.total_size; + + const auto& info = image.info; + const u32 bytes_per_block = BytesPerBlock(info.format); + const u32 width_blocks = Common::DivCeil(info.size.width, 4u); + const u32 height_blocks = Common::DivCeil(info.size.height, 4u); + task.bytes_per_slice = static_cast(width_blocks * bytes_per_block) * height_blocks; + task.last_submitted_offset = 0; + + const auto sp = VideoCommon::Accelerated::MakeBlockLinearSwizzle3DParams( + uploads[0], task.info); + task.swizzled_slice_size = sp.slice_size; + task.swizzle_block_depth = sp.block_depth; + + const u32 total_slices_init = info.size.depth; + + task.is_sparse = True(image.flags & ImageFlagBits::Sparse); + task.active_z_start = 0; + task.active_z_end = total_slices_init; + if (task.is_sparse) { + const auto segs = + gpu_memory->GetSubmappedRange(image.gpu_addr, image.guest_size_bytes); + task.sparse_segments.assign(segs.begin(), segs.end()); + task.segment_scan_cursor = 0; + + task.slice_has_data.assign(image.info.size.depth, 0u); + + if (image.info.size.depth > 1 && !image.slice_offsets.empty()) { + const u32 depth = static_cast(image.info.size.depth); + const u32 group_slices_sparse = 1u << task.swizzle_block_depth; + const u64 group_byte_size = task.swizzled_slice_size; + + size_t seg_idx = 0; + for (u32 g_start = 0; g_start < depth; g_start += group_slices_sparse) { + const u32 g_end = (std::min)(g_start + group_slices_sparse, depth); + const u64 group_abs_start = image.slice_offsets[g_start]; + const u64 group_abs_end = + (g_end < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[g_end]) + : group_abs_start + group_byte_size; + + while (seg_idx < task.sparse_segments.size()) { + const auto& [seg_gpu_addr, seg_size] = task.sparse_segments[seg_idx]; + const u64 seg_start = seg_gpu_addr - image.gpu_addr; + const u64 seg_end = seg_start + seg_size; + if (seg_end <= group_abs_start) { + ++seg_idx; + continue; + } + break; + } + + bool group_has_data = false; + for (size_t probe = seg_idx; probe < task.sparse_segments.size(); ++probe) { + const auto& [seg_gpu_addr, seg_size] = task.sparse_segments[probe]; + const u64 seg_start = seg_gpu_addr - image.gpu_addr; + const u64 seg_end = seg_start + seg_size; + if (seg_start >= group_abs_end) break; + if (seg_end > group_abs_start && seg_start < group_abs_end) { + group_has_data = true; + break; + } + } + + if (group_has_data) { + for (u32 z = g_start; z < g_end; ++z) { + task.slice_has_data[z] = 1u; + } + } + } + } else { + std::fill(task.slice_has_data.begin(), task.slice_has_data.end(), 1u); + } + + if (async_unswizzle_slices_per_batch == 0) { + u32 first_populated = total_slices_init; + u32 last_populated_excl = 0; + for (u32 z = 0; z < total_slices_init; ++z) { + if (task.slice_has_data[z]) { + first_populated = (std::min)(first_populated, z); + last_populated_excl = z + 1; + } + } + if (first_populated < last_populated_excl) { + const u32 group_slices_align = (std::max)(1u << task.swizzle_block_depth, 1u); + task.active_z_start = (first_populated / group_slices_align) * group_slices_align; + task.active_z_end = (std::min)( + Common::AlignUp(last_populated_excl, group_slices_align), total_slices_init); + } else { + task.active_z_start = 0; + task.active_z_end = 0; + } + } + } + + const u32 batch_slices = (async_unswizzle_slices_per_batch == 0) + ? (task.active_z_end - task.active_z_start) + : async_unswizzle_slices_per_batch; + + size_t calculated_max = 0; + const u32 group_slices = (std::min)(1u << task.swizzle_block_depth, total_slices_init); + + // Complicated but I'd rather size the staging buffer to the slices we found rather than the whole texture + // Not every machine has 32GB of ram + for (u32 z = task.active_z_start; z < task.active_z_end; z += (std::max)(batch_slices, 1u)) { + const u32 z_want_init = (std::min)(z + (std::max)(batch_slices, 1u), task.active_z_end); + const u32 z_count_init = z_want_init - z; + + const u32 z_group_start = (z / group_slices) * group_slices; + const u32 z_group_end = Common::AlignUp(z + z_count_init, group_slices); + const u32 z_group_count = (std::min)(z_group_end, total_slices_init) - z_group_start; + + const size_t group_off = (z_group_start < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start]) : 0; + const size_t group_end = (z_group_start + z_group_count < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start + z_group_count]) : task.total_size; + + calculated_max = (std::max)(calculated_max, group_end - group_off); + } + if (calculated_max > 0 && calculated_max < max_batch_size) { + max_batch_size = calculated_max; + } + if (task.active_z_start >= task.active_z_end) { + max_batch_size = (std::min)(max_batch_size, task.bytes_per_slice); + } + + if (task.active_z_start > 0 && task.active_z_start < image.slice_offsets.size()) { + const size_t start_byte = static_cast(image.slice_offsets[task.active_z_start]); + task.current_offset = start_byte; + task.last_submitted_offset = task.active_z_start; + } + + // Possible memory leak here but if UnswizzleSharedStagingCap does its job we should be okay + const size_t needed = (std::max)(max_batch_size, size_t{1}); + if (needed > UnswizzleSharedStagingCap) { + task.staging_buffer = runtime.UploadStagingBuffer(needed, true); + task.owns_staging_buffer = true; + } else { + if (!unswizzle_shared_staging.has_value() || + unswizzle_shared_staging_capacity < needed) { + if (unswizzle_shared_staging.has_value()) { + runtime.FreeDeferredStagingBuffer(*unswizzle_shared_staging); + } + const size_t grown = (std::min)( + UnswizzleSharedStagingCap, + (std::max)(needed, unswizzle_shared_staging_capacity + + unswizzle_shared_staging_capacity / 4)); + unswizzle_shared_staging = runtime.UploadStagingBuffer(grown, true); + unswizzle_shared_staging_capacity = grown; + } + task.staging_buffer = *unswizzle_shared_staging; + task.staging_buffer.mapped_span = + task.staging_buffer.mapped_span.subspan(0, needed); + task.owns_staging_buffer = false; + } + + task.initialized = true; +} + +template +void TextureCache

::StageSwizzledDataCpu(PendingUnswizzle& task, Image& image, + size_t current_group_off, size_t current_group_end) { + if (!task.cpu_job_in_flight && task.current_offset < current_group_end) { + u8* const staging_base = task.staging_buffer.mapped_span.data(); + const size_t group_window_size = current_group_end - current_group_off; + + if (task.current_offset == current_group_off) { + std::memset(staging_base, 0, (std::min)(group_window_size, task.staging_buffer.mapped_span.size())); + } + + const size_t remaining = current_group_end - task.current_offset; + size_t copy_amount = (async_unswizzle_chunk_size == 0) + ? remaining + : (std::min)(async_unswizzle_chunk_size, remaining); + + if (async_unswizzle_chunk_size > 0 && copy_amount < remaining) { + copy_amount = (copy_amount / task.bytes_per_slice) * task.bytes_per_slice; + if (copy_amount == 0) copy_amount = task.bytes_per_slice; + copy_amount = (std::min)(copy_amount, remaining); + } + + const size_t base_off = task.staging_base_byte_offset; + const size_t read_start = task.current_offset; + const size_t read_end = read_start + copy_amount; + + const size_t buffer_offset_start = read_start - current_group_off; + const size_t abs_start = read_start + base_off; + const size_t abs_end = read_end + base_off; + + if (task.is_sparse) { + const size_t max_gap_for_batch = 4096; // This needs tuning as its just a number I pulled from out of nowhere + + while (task.segment_scan_cursor < task.sparse_segments.size()) { + const auto& [seg_gpu_addr, seg_size] = + task.sparse_segments[task.segment_scan_cursor]; + const size_t seg_abs_start = static_cast(seg_gpu_addr - image.gpu_addr); + const size_t seg_abs_end = seg_abs_start + seg_size; + + if (seg_abs_end <= abs_start) { ++task.segment_scan_cursor; continue; } + if (seg_abs_start >= abs_end) { break; } + + const size_t ol_abs_start = (std::max)(seg_abs_start, abs_start); + const size_t ol_abs_end = (std::min)(seg_abs_end, abs_end); + + size_t batch_abs_end = ol_abs_end; + size_t batch_cursor = task.segment_scan_cursor + 1; + + while (batch_cursor < task.sparse_segments.size()) { + const auto& [next_seg_addr, next_seg_size] = task.sparse_segments[batch_cursor]; + const size_t next_seg_abs_start = static_cast(next_seg_addr - image.gpu_addr); + const size_t next_seg_abs_end = next_seg_abs_start + next_seg_size; + + if (next_seg_abs_start >= abs_end) break; + if (next_seg_abs_end <= abs_start) { ++batch_cursor; continue; } + + const size_t gap = next_seg_abs_start - batch_abs_end; + if (gap > max_gap_for_batch) break; + + batch_abs_end = (std::min)(next_seg_abs_end, abs_end); + ++batch_cursor; + } + + const size_t batch_size = batch_abs_end - ol_abs_start; + const size_t buffer_ol_rel_start = (ol_abs_start - base_off) - current_group_off; + + gpu_memory->ReadBlockUnsafe(image.gpu_addr + ol_abs_start, + staging_base + buffer_ol_rel_start, + batch_size); + + if (seg_abs_end > abs_end) break; + task.segment_scan_cursor = batch_cursor; + } + } else { + gpu_memory->ReadBlockUnsafe(image.gpu_addr + abs_start, staging_base + buffer_offset_start, + copy_amount); + } + task.current_offset += copy_amount; + } +} + +template +void TextureCache

::DispatchCpuUnswizzleJob(PendingUnswizzle& task, Image& image, u32 z_src, u32 z_count) { + const u32 total_slices = image.info.size.depth; + const u32 z_image = z_src; + + if (z_count > 0) { + bool any_data = true; + if (task.is_sparse && !task.slice_has_data.empty()) { + any_data = false; + for (u32 z = z_src; z < z_src + z_count; ++z) { + if (z < task.slice_has_data.size() && task.slice_has_data[z]) { + any_data = true; + break; + } + } + } + + if (!any_data) { + // I had no choice but to do this otherwise the area around the data within the texture will be filled with uninitialized junk from VRAM + const size_t upload_size = static_cast(z_count) * task.bytes_per_slice; + auto upload_staging = runtime.UploadStagingBuffer(upload_size); + std::memset(upload_staging.mapped_span.data(), 0, upload_size); + + const u32 aligned_width = Common::AlignUp(image.info.size.width, 4u); + const u32 aligned_height = Common::AlignUp(image.info.size.height, 4u); + + const BufferImageCopy copy{ + .buffer_offset = 0, + .buffer_size = upload_size, + .buffer_row_length = aligned_width, + .buffer_image_height = aligned_height, + .image_subresource = + { + .base_level = 0, + .base_layer = 0, + .num_layers = 1, + }, + .image_offset = {0, 0, static_cast(z_image)}, + .image_extent = {image.info.size.width, image.info.size.height, z_count}, + }; + const std::array copies{copy}; + image.UploadMemory(upload_staging, copies); + runtime.InsertUploadMemoryBarrier(); + + task.last_submitted_offset += z_count; + } else { + const u32 group_slices = (std::min)(1u << task.swizzle_block_depth, total_slices); + const u32 z_group_start = (z_src / group_slices) * group_slices; + const u32 z_group_end_excl = Common::AlignUp(z_src + z_count, group_slices); + const u32 z_group_count = (std::min)(z_group_end_excl, total_slices) - z_group_start; + + task.cpu_job_in_flight = true; + if (!task.cpu_chunk) { + task.cpu_chunk = std::make_unique(); + } + auto* chunk = task.cpu_chunk.get(); + chunk->complete = false; + chunk->z_src = z_src; + chunk->z_image = z_image; + chunk->z_count = z_count; + + const size_t group_guest_off = (z_group_start < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start]) : 0; + const size_t group_guest_end = (z_group_start + z_group_count < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start + z_group_count]) : task.total_size; + const size_t group_guest_size = group_guest_end - group_guest_off; + + const u8* const staging_base = task.staging_buffer.mapped_span.data(); + chunk->swizzled_batch = std::span(staging_base, group_guest_size); + + const size_t group_linear_size = static_cast(z_group_count) * task.bytes_per_slice; + if (chunk->linear_batch.size() < group_linear_size) { + chunk->linear_batch.resize(group_linear_size); + } + std::memset(chunk->linear_batch.data(), 0, group_linear_size); + + chunk->group_z_start = z_group_start; + chunk->group_z_count = z_group_count; + + boost::container::small_vector, 32> subranges; + for (u32 sub_start = 0; sub_start < z_group_count; sub_start += group_slices) { + const u32 sub_count = (std::min)(group_slices, z_group_count - sub_start); + subranges.emplace_back(sub_start, sub_count); + } + + chunk->subjobs_pending = static_cast(subranges.size()); + + for (const auto& [sub_start, sub_count] : subranges) { + const size_t sub_guest_off = (z_group_start + sub_start < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start + sub_start]) - group_guest_off + : (static_cast(sub_start) * task.bytes_per_slice); + const size_t sub_guest_end = + (z_group_start + sub_start + sub_count < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start + sub_start + sub_count]) - + group_guest_off + : group_guest_size; + const size_t sub_guest_size = sub_guest_end - sub_guest_off; + + const size_t sub_linear_off = static_cast(sub_start) * task.bytes_per_slice; + const size_t sub_linear_size = static_cast(sub_count) * task.bytes_per_slice; + + std::span sub_swizzled = + chunk->swizzled_batch.subspan(sub_guest_off, sub_guest_size); + std::span sub_linear(chunk->linear_batch.data() + sub_linear_off, + sub_linear_size); + + auto func = [chunk, sub_linear, sub_swizzled, sub_count, + num_tiles = task.cpu_num_tiles, block = task.cpu_block, + stride_alignment = task.cpu_stride_alignment, + bytes_per_block = task.cpu_bytes_per_block]() mutable { + UnswizzleTexture(sub_linear, sub_swizzled, bytes_per_block, + num_tiles.width, num_tiles.height, sub_count, + block.height, block.depth, stride_alignment); + if (chunk->subjobs_pending.fetch_sub(1, std::memory_order_acq_rel) == 1) { + std::unique_lock lock{chunk->mutex}; + chunk->complete = true; + } + }; + texture_decode_worker.QueueWork(std::move(func)); + } + } + } +} + +template +void TextureCache

::PollAndUploadCpuUnswizzleJob(PendingUnswizzle& task, Image& image) { + if (task.cpu_job_in_flight) { + auto& chunk = *task.cpu_chunk; + { + std::unique_lock lock{chunk.mutex}; + if (!chunk.complete) { + return; + } + } + + const size_t offset_slices = chunk.z_src - chunk.group_z_start; + const size_t upload_offset = offset_slices * task.bytes_per_slice; + const size_t upload_size = static_cast(chunk.z_count) * task.bytes_per_slice; + + auto upload_staging = runtime.UploadStagingBuffer(upload_size); + std::memcpy(upload_staging.mapped_span.data(), chunk.linear_batch.data() + upload_offset, + upload_size); + + const u32 aligned_width = Common::AlignUp(image.info.size.width, 4u); + const u32 aligned_height = Common::AlignUp(image.info.size.height, 4u); + + const BufferImageCopy copy{ + .buffer_offset = 0, + .buffer_size = upload_size, + .buffer_row_length = aligned_width, + .buffer_image_height = aligned_height, + .image_subresource = + { + .base_level = 0, + .base_layer = 0, + .num_layers = 1, + }, + .image_offset = {0, 0, static_cast(chunk.z_image)}, + .image_extent = {image.info.size.width, image.info.size.height, chunk.z_count}, + }; + const std::array copies{copy}; + image.UploadMemory(upload_staging, copies); + runtime.InsertUploadMemoryBarrier(); + + task.last_submitted_offset += chunk.z_count; + task.cpu_job_in_flight = false; + } +} + +template +void TextureCache

::TickAsyncUnswizzleCpu(PendingUnswizzle& task, Image& image) { + if (!task.initialized) { + InitializeCpuUnswizzleTask(task, image); + } + + const u32 total_slices = image.info.size.depth; + const u32 effective_end = (async_unswizzle_slices_per_batch == 0 && task.is_sparse) + ? task.active_z_end + : total_slices; + const u32 batch = (async_unswizzle_slices_per_batch == 0) + ? effective_end + : async_unswizzle_slices_per_batch; + + const u32 z_src = static_cast(task.last_submitted_offset); + const u32 z_want = (std::min)(z_src + batch, effective_end); + const u32 z_count_current = (z_want > z_src) ? (z_want - z_src) : 0; + + size_t current_group_off = 0; + size_t current_group_end = task.total_size; + + if (z_count_current > 0) { + const u32 group_slices = (std::min)(1u << task.swizzle_block_depth, total_slices); + const u32 z_group_start = (z_src / group_slices) * group_slices; + const u32 z_group_end_excl = Common::AlignUp(z_src + z_count_current, group_slices); + const u32 z_group_count = (std::min)(z_group_end_excl, total_slices) - z_group_start; + + current_group_off = (z_group_start < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start]) : 0; + current_group_end = (z_group_start + z_group_count < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[z_group_start + z_group_count]) : task.total_size; + } else { + current_group_off = task.total_size; + current_group_end = task.total_size; + } + + StageSwizzledDataCpu(task, image, current_group_off, current_group_end); + + const size_t active_end_byte = (task.active_z_end < image.slice_offsets.size()) + ? static_cast(image.slice_offsets[task.active_z_end]) : task.total_size; + const bool is_final_batch = task.current_offset >= active_end_byte; + const bool batch_ready = task.current_offset >= current_group_end; + + if (!task.cpu_job_in_flight && z_src < effective_end && + (batch_ready || is_final_batch)) { + DispatchCpuUnswizzleJob(task, image, z_src, z_count_current); + } + + PollAndUploadCpuUnswizzleJob(task, image); + + // Check if complete + const u32 slices_submitted = static_cast(task.last_submitted_offset); + const bool all_submitted = slices_submitted >= effective_end; + + if (is_final_batch && all_submitted) { + if (task.owns_staging_buffer) { + runtime.FreeDeferredStagingBuffer(task.staging_buffer); + } + image.flags &= ~ImageFlagBits::IsDecoding; + unswizzle_queue.pop_front(); + } +} + template bool TextureCache

::ScaleUp(Image& image) { const bool has_copy = image.HasScaled(); @@ -2479,6 +3091,26 @@ void TextureCache

::UntrackImage(ImageBase& image, ImageId image_id) { template void TextureCache

::DeleteImage(ImageId image_id, bool immediate_delete) { ImageBase& image = slot_images[image_id]; + if (True(image.flags & ImageFlagBits::IsDecoding)) { + const bool has_cpu_job_in_flight = + std::ranges::any_of(unswizzle_queue, [image_id](const PendingUnswizzle& task) { + return task.image_id == image_id && task.cpu_job_in_flight; + }); + if (has_cpu_job_in_flight) { + texture_decode_worker.WaitForRequests(); + } + for (auto& task : unswizzle_queue) { + if (task.image_id == image_id && task.initialized && task.owns_staging_buffer) { + runtime.FreeDeferredStagingBuffer(task.staging_buffer); + } + } + std::erase_if(unswizzle_queue, + [image_id](const PendingUnswizzle& task) { return task.image_id == image_id; }); + std::erase_if(async_decodes, [image_id](const std::unique_ptr& ctx) { + return ctx->image_id == image_id; + }); + image.flags &= ~ImageFlagBits::IsDecoding; + } if (image.HasScaled()) { total_used_memory -= GetScaledImageSizeBytes(image); } diff --git a/src/video_core/texture_cache/texture_cache_base.h b/src/video_core/texture_cache/texture_cache_base.h index 03b11298ba..ccc39a2f74 100644 --- a/src/video_core/texture_cache/texture_cache_base.h +++ b/src/video_core/texture_cache/texture_cache_base.h @@ -10,7 +10,9 @@ #include #include #include +#include #include +#include #include // TODO: find out which don't require stable iters #include @@ -70,6 +72,19 @@ struct AsyncDecodeContext { std::atomic_bool complete; }; +struct AsyncCpuUnswizzleChunk { + std::span swizzled_batch; + std::vector linear_batch; + u32 z_src = 0; + u32 z_image = 0; + u32 z_count = 0; + u32 group_z_start = 0; + u32 group_z_count = 0; + std::mutex mutex; + std::atomic_bool complete{false}; + std::atomic subjobs_pending{0}; +}; + using TextureCacheGPUMap = ankerl::unordered_dense::map, Common::IdentityHash>; class TextureCacheChannelInfo : public ChannelInfo { @@ -132,26 +147,38 @@ class TextureCache : public VideoCommon::ChannelSetupCaches slice_has_data; + size_t staging_base_byte_offset = 0; + size_t bytes_per_slice = 0; + u64 swizzled_slice_size = 0; + std::vector> sparse_segments; + std::vector slice_has_data; size_t segment_scan_cursor = 0; - size_t swizzle_group_size = 0; - u32 slices_per_group = 0; - bool is_incremental = false; - size_t staging_base_byte_offset = 0; - u32 incremental_z_start = 0; - u32 incremental_z_count = 0; + u32 active_z_start = 0; + u32 active_z_end = 0; + + std::unique_ptr cpu_chunk; + Extent3D cpu_num_tiles{}; + Extent3D cpu_block{}; + + u32 swizzle_block_depth = 0; + u32 cpu_stride_alignment = 0; + u32 cpu_bytes_per_block = 0; + + bool initialized = false; + bool is_sparse = false; + bool is_cpu = false; + bool cpu_job_in_flight = false; }; struct BlitImages { @@ -438,6 +465,13 @@ private: void QueueAsyncUnswizzle(Image& image, ImageId image_id); void TickAsyncUnswizzle(); + void TickAsyncUnswizzleGpu(PendingUnswizzle& task, Image& image); + void TickAsyncUnswizzleCpu(PendingUnswizzle& task, Image& image); + void InitializeCpuUnswizzleTask(PendingUnswizzle& task, Image& image); + void StageSwizzledDataCpu(PendingUnswizzle& task, Image& image, size_t current_group_off, + size_t current_group_end); + void DispatchCpuUnswizzleJob(PendingUnswizzle& task, Image& image, u32 z_src, u32 z_count); + void PollAndUploadCpuUnswizzleJob(PendingUnswizzle& task, Image& image); bool IsUnswizzleStorageFormatSupported(PixelFormat format) { return runtime.IsUnswizzleStorageFormatSupported(format); @@ -472,9 +506,11 @@ private: u64 minimum_memory; u64 expected_memory; u64 critical_memory; - size_t gpu_unswizzle_maxsize = 0; - size_t swizzle_chunk_size = 0; - u32 swizzle_slices_per_batch = 0; + + Settings::AsyncUnswizzleMode async_unswizzle_mode = Settings::AsyncUnswizzleMode::Off; + size_t async_unswizzle_maxsize = 0; + size_t async_unswizzle_chunk_size = 0; + u32 async_unswizzle_slices_per_batch = 0; struct BufferDownload { GPUVAddr address; @@ -527,12 +563,25 @@ private: u64 frame_tick = 0; u64 last_sampler_gc_frame = (std::numeric_limits::max)(); + Common::ThreadWorker texture_decode_worker{1, "TextureDecoder"}; + // I kinda don't want ASTC CPU async to flood your threads but eh, lets FAFO + static u32 ComputeTextureDecodeWorkerCount() { + const u32 hw = std::thread::hardware_concurrency(); + return (std::max)(1u, hw > 2 ? hw - 1 : hw); + } + const u32 texture_decode_worker_count = ComputeTextureDecodeWorkerCount(); + Common::ThreadWorker texture_decode_worker{texture_decode_worker_count, "TextureDecoder"}; Common::ThreadWorker texture_decode_worker{1, "TextureDecoder", {}, Common::ThreadPlacement::Efficiency}; std::vector> async_decodes; std::deque unswizzle_queue; + static constexpr size_t UnswizzleSharedStagingCap = 1_GiB; + std::optional unswizzle_shared_staging; + size_t unswizzle_shared_staging_capacity = 0; + bool unswizzle_shared_staging_pending_gpu_read = false; + // Join caching boost::container::small_vector join_overlap_ids; ankerl::unordered_dense::set join_overlaps_found; diff --git a/src/video_core/textures/decoders.cpp b/src/video_core/textures/decoders.cpp index 25b0860b0f..efb2e79cef 100644 --- a/src/video_core/textures/decoders.cpp +++ b/src/video_core/textures/decoders.cpp @@ -1,4 +1,4 @@ -// SPDX-FileCopyrightText: Copyright 2025 Eden Emulator Project +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project // SPDX-License-Identifier: GPL-3.0-or-later // SPDX-FileCopyrightText: Copyright 2018 yuzu Emulator Project @@ -12,9 +12,14 @@ #include "common/alignment.h" #include "common/assert.h" #include "common/bit_util.h" +#include "common/cpu_features.h" #include "common/div_ceil.h" #include "video_core/gpu.h" #include "video_core/textures/decoders.h" +#if defined(ARCHITECTURE_x86_64) +#include "video_core/textures/decoders_avx2.h" +#include "video_core/textures/decoders_sse2.h" +#endif namespace Tegra::Texture { namespace { @@ -58,34 +63,90 @@ void SwizzleImpl(std::span output, std::span input, u32 width, u32 const u32 block_depth_mask = (1U << block_depth) - 1; const u32 x_shift = GOB_SIZE_SHIFT + block_height + block_depth; - for (u32 slice = 0; slice < depth; ++slice) { - const u32 z = slice + origin_z; - const u32 offset_z = (z >> block_depth) * slice_size + - ((z & block_depth_mask) << (GOB_SIZE_SHIFT + block_height)); - for (u32 line = 0; line < height; ++line) { - const u32 y = line + origin_y; - const u32 swizzled_y = pdep(y); + if constexpr (BYTES_PER_PIXEL == 1 || BYTES_PER_PIXEL == 2 || BYTES_PER_PIXEL == 4 || + BYTES_PER_PIXEL == 8 || BYTES_PER_PIXEL == 16) { + static constexpr u32 RUN_BYTES = 16; + static constexpr u32 PIXELS_PER_RUN = RUN_BYTES / BYTES_PER_PIXEL; - const u32 block_y = y >> GOB_SIZE_Y_SHIFT; - const u32 offset_y = (block_y >> block_height) * block_size + - ((block_y & block_height_mask) << GOB_SIZE_SHIFT); + for (u32 slice = 0; slice < depth; ++slice) { + const u32 z = slice + origin_z; + const u32 offset_z = (z >> block_depth) * slice_size + + ((z & block_depth_mask) << (GOB_SIZE_SHIFT + block_height)); + for (u32 line = 0; line < height; ++line) { + const u32 y = line + origin_y; + const u32 swizzled_y = pdep(y); - u32 swizzled_x = pdep(origin_x * BYTES_PER_PIXEL); - for (u32 column = 0; column < width; - ++column, incrpdep(swizzled_x)) { - const u32 x = (column + origin_x) * BYTES_PER_PIXEL; - const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 block_y = y >> GOB_SIZE_Y_SHIFT; + const u32 offset_y = (block_y >> block_height) * block_size + + ((block_y & block_height_mask) << GOB_SIZE_SHIFT); - const u32 base_swizzled_offset = offset_z + offset_y + offset_x; - const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + u32 swizzled_x = pdep(origin_x * BYTES_PER_PIXEL); + u32 column = 0; - const u32 unswizzled_offset = - slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + for (; column + PIXELS_PER_RUN <= width; + column += PIXELS_PER_RUN, incrpdep(swizzled_x)) { + const u32 x = (column + origin_x) * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; - u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset]; - const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset]; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); - std::memcpy(dst, src, BYTES_PER_PIXEL); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset]; + const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset]; + + std::memcpy(dst, src, RUN_BYTES); + } + for (; column < width; + ++column, incrpdep(swizzled_x)) { + const u32 x = (column + origin_x) * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset]; + const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset]; + + std::memcpy(dst, src, BYTES_PER_PIXEL); + } + } + } + } else { + for (u32 slice = 0; slice < depth; ++slice) { + const u32 z = slice + origin_z; + const u32 offset_z = (z >> block_depth) * slice_size + + ((z & block_depth_mask) << (GOB_SIZE_SHIFT + block_height)); + for (u32 line = 0; line < height; ++line) { + const u32 y = line + origin_y; + const u32 swizzled_y = pdep(y); + + const u32 block_y = y >> GOB_SIZE_Y_SHIFT; + const u32 offset_y = (block_y >> block_height) * block_size + + ((block_y & block_height_mask) << GOB_SIZE_SHIFT); + + u32 swizzled_x = pdep(origin_x * BYTES_PER_PIXEL); + for (u32 column = 0; column < width; + ++column, incrpdep(swizzled_x)) { + const u32 x = (column + origin_x) * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[TO_LINEAR ? swizzled_offset : unswizzled_offset]; + const u8* const src = &input[TO_LINEAR ? unswizzled_offset : swizzled_offset]; + + std::memcpy(dst, src, BYTES_PER_PIXEL); + } } } } @@ -157,6 +218,32 @@ void SwizzleSubrectImpl(std::span output, std::span input, u32 wid template void Swizzle(std::span output, std::span input, u32 bytes_per_pixel, u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, u32 stride_alignment) { +// This is quite ugly and has duplicate code but I want to make sure this is an actual speedup on stuff like SteamDeck +#if defined(ARCHITECTURE_x86_64) + if constexpr (!TO_LINEAR) { + switch (bytes_per_pixel) { + case 1: + case 2: + case 4: + case 8: + case 16: + // This needs more testing but in my benchmarks this is roughly 1.5x faster than master + // A check for AVX-512 could be possible but my CPU doesn't support it so no way to test it + if (Common::g_cpu_caps.avx2) { + UnswizzleGobPermuteAVX2(output, input, bytes_per_pixel, width, height, depth, + block_height, block_depth, stride_alignment); + } else { + // SSE2 comes with x64 by default but incase I'm wrong this is here + // Could maybe check for SSE3, SSSE3, SSE4.1 and SSE4.2 but I think they don't really add anything for this scenario + UnswizzleGobPermuteSSE2(output, input, bytes_per_pixel, width, height, depth, + block_height, block_depth, stride_alignment); + } + return; + default: + break; + } + } +#endif switch (bytes_per_pixel) { #define BPP_CASE(x) \ case x: \ diff --git a/src/video_core/textures/decoders_avx2.cpp b/src/video_core/textures/decoders_avx2.cpp new file mode 100644 index 0000000000..56604d3865 --- /dev/null +++ b/src/video_core/textures/decoders_avx2.cpp @@ -0,0 +1,157 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later + +#include +#include + +#include // AVX2 + +#include "common/assert.h" +#include "common/div_ceil.h" +#include "video_core/textures/decoders.h" +#include "video_core/textures/decoders_avx2.h" + +namespace Tegra::Texture { +namespace { + +template +constexpr u32 Pdep(u32 value) { + u32 result = 0; + u32 m = mask; + for (u32 bit = 1; m; bit += bit) { + if (value & bit) { + result |= m & (~m + 1); + } + m &= m - 1; + } + return result; +} + +template +void IncrPdep(u32& value) { + static constexpr u32 swizzled_incr = Pdep(incr_amount); + value = ((value | ~mask) + swizzled_incr) & mask; +} + +template +__attribute__((target("avx2"))) void UnswizzleGobPermuteAVX2Impl(std::span output, + std::span input, + u32 width, u32 height, u32 depth, + u32 block_height, + u32 block_depth, u32 stride) { + static_assert(32 % BYTES_PER_PIXEL == 0, "BYTES_PER_PIXEL must evenly divide 32"); + + static constexpr u32 RUN_BYTES = 32; + static constexpr u32 PIXELS_PER_RUN = RUN_BYTES / BYTES_PER_PIXEL; + + const u32 pitch = width * BYTES_PER_PIXEL; + const u32 gobs_in_x = Common::DivCeilLog2(stride, GOB_SIZE_X_SHIFT); + const u32 block_size = gobs_in_x << (GOB_SIZE_SHIFT + block_height + block_depth); + const u32 slice_size = + Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size; + const u32 block_height_mask = (1U << block_height) - 1; + const u32 block_depth_mask = (1U << block_depth) - 1; + const u32 x_shift = GOB_SIZE_SHIFT + block_height + block_depth; + + for (u32 slice = 0; slice < depth; ++slice) { + const u32 offset_z = (slice >> block_depth) * slice_size + + ((slice & block_depth_mask) << (GOB_SIZE_SHIFT + block_height)); + + for (u32 line = 0; line < height; ++line) { + const u32 swizzled_y = Pdep(line); + const u32 block_y = line >> GOB_SIZE_Y_SHIFT; + const u32 offset_y = (block_y >> block_height) * block_size + + ((block_y & block_height_mask) << GOB_SIZE_SHIFT); + + u32 swizzled_x = 0; + u32 column = 0; + + for (; column + PIXELS_PER_RUN <= width; + column += PIXELS_PER_RUN, IncrPdep(swizzled_x)) { + + const u32 x = column * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[unswizzled_offset]; + const u8* const src = &input[swizzled_offset]; + + const __m128i lo = _mm_loadu_si128(reinterpret_cast(src)); + const __m128i hi = _mm_loadu_si128(reinterpret_cast(src + 32)); + + __m256i val = _mm256_castsi128_si256(lo); + val = _mm256_inserti128_si256(val, hi, 1); + + _mm256_storeu_si256(reinterpret_cast<__m256i*>(dst), val); + } + + if constexpr (BYTES_PER_PIXEL <= 16) { + constexpr u32 PIXELS_PER_16 = 16 / BYTES_PER_PIXEL; + if (column + PIXELS_PER_16 <= width) { + const u32 x = column * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[unswizzled_offset]; + const u8* const src = &input[swizzled_offset]; + + const __m128i val = _mm_loadu_si128(reinterpret_cast(src)); + _mm_storeu_si128(reinterpret_cast<__m128i*>(dst), val); + + column += PIXELS_PER_16; + IncrPdep(swizzled_x); + } + } + + for (; column < width; + ++column, IncrPdep(swizzled_x)) { + + const u32 x = column * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[unswizzled_offset]; + const u8* const src = &input[swizzled_offset]; + + std::memcpy(dst, src, BYTES_PER_PIXEL); + } + } + } +} + +} // Anonymous namespace + +void UnswizzleGobPermuteAVX2(std::span output, std::span input, u32 bytes_per_pixel, + u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, + u32 stride) { + switch (bytes_per_pixel) { + case 1: + return UnswizzleGobPermuteAVX2Impl<1>(output, input, width, height, depth, block_height, + block_depth, stride); + case 2: + return UnswizzleGobPermuteAVX2Impl<2>(output, input, width, height, depth, block_height, + block_depth, stride); + case 4: + return UnswizzleGobPermuteAVX2Impl<4>(output, input, width, height, depth, block_height, + block_depth, stride); + case 8: + return UnswizzleGobPermuteAVX2Impl<8>(output, input, width, height, depth, block_height, + block_depth, stride); + case 16: + return UnswizzleGobPermuteAVX2Impl<16>(output, input, width, height, depth, block_height, + block_depth, stride); + default: + UNREACHABLE(); + } +} + +} // namespace Tegra::Texture \ No newline at end of file diff --git a/src/video_core/textures/decoders_avx2.h b/src/video_core/textures/decoders_avx2.h new file mode 100644 index 0000000000..cb113ffa93 --- /dev/null +++ b/src/video_core/textures/decoders_avx2.h @@ -0,0 +1,16 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later + +#pragma once + +#include + +#include "common/common_types.h" + +namespace Tegra::Texture { + +void UnswizzleGobPermuteAVX2(std::span output, std::span input, u32 bytes_per_pixel, + u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, + u32 stride); + +} // namespace Tegra::Texture diff --git a/src/video_core/textures/decoders_sse2.cpp b/src/video_core/textures/decoders_sse2.cpp new file mode 100644 index 0000000000..071c0df039 --- /dev/null +++ b/src/video_core/textures/decoders_sse2.cpp @@ -0,0 +1,131 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later + +#include +#include + +#include // SSE2 + +#include "common/assert.h" +#include "common/div_ceil.h" +#include "video_core/textures/decoders.h" +#include "video_core/textures/decoders_sse2.h" + +namespace Tegra::Texture { +namespace { + +template +constexpr u32 Pdep(u32 value) { + u32 result = 0; + u32 m = mask; + for (u32 bit = 1; m; bit += bit) { + if (value & bit) { + result |= m & (~m + 1); + } + m &= m - 1; + } + return result; +} + +template +void IncrPdep(u32& value) { + static constexpr u32 swizzled_incr = Pdep(incr_amount); + value = ((value | ~mask) + swizzled_incr) & mask; +} + +template +__attribute__((target("sse2"))) void UnswizzleGobPermuteSSE2Impl(std::span output, + std::span input, + u32 width, u32 height, u32 depth, + u32 block_height, + u32 block_depth, u32 stride) { + static_assert(16 % BYTES_PER_PIXEL == 0, "BYTES_PER_PIXEL must evenly divide 16"); + + static constexpr u32 RUN_BYTES = 16; + static constexpr u32 PIXELS_PER_RUN = RUN_BYTES / BYTES_PER_PIXEL; + + const u32 pitch = width * BYTES_PER_PIXEL; + const u32 gobs_in_x = Common::DivCeilLog2(stride, GOB_SIZE_X_SHIFT); + const u32 block_size = gobs_in_x << (GOB_SIZE_SHIFT + block_height + block_depth); + const u32 slice_size = + Common::DivCeilLog2(height, block_height + GOB_SIZE_Y_SHIFT) * block_size; + const u32 block_height_mask = (1U << block_height) - 1; + const u32 block_depth_mask = (1U << block_depth) - 1; + const u32 x_shift = GOB_SIZE_SHIFT + block_height + block_depth; + + for (u32 slice = 0; slice < depth; ++slice) { + const u32 offset_z = (slice >> block_depth) * slice_size + + ((slice & block_depth_mask) << (GOB_SIZE_SHIFT + block_height)); + + for (u32 line = 0; line < height; ++line) { + const u32 swizzled_y = Pdep(line); + const u32 block_y = line >> GOB_SIZE_Y_SHIFT; + const u32 offset_y = (block_y >> block_height) * block_size + + ((block_y & block_height_mask) << GOB_SIZE_SHIFT); + + u32 swizzled_x = 0; + u32 column = 0; + + for (; column + PIXELS_PER_RUN <= width; + column += PIXELS_PER_RUN, IncrPdep(swizzled_x)) { + + const u32 x = column * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[unswizzled_offset]; + const u8* const src = &input[swizzled_offset]; + + const __m128i val = _mm_loadu_si128(reinterpret_cast(src)); + _mm_storeu_si128(reinterpret_cast<__m128i*>(dst), val); + } + + for (; column < width; + ++column, IncrPdep(swizzled_x)) { + + const u32 x = column * BYTES_PER_PIXEL; + const u32 offset_x = (x >> GOB_SIZE_X_SHIFT) << x_shift; + const u32 base_swizzled_offset = offset_z + offset_y + offset_x; + const u32 swizzled_offset = base_swizzled_offset + (swizzled_x | swizzled_y); + const u32 unswizzled_offset = + slice * pitch * height + line * pitch + column * BYTES_PER_PIXEL; + + u8* const dst = &output[unswizzled_offset]; + const u8* const src = &input[swizzled_offset]; + + std::memcpy(dst, src, BYTES_PER_PIXEL); + } + } + } +} + +} // Anonymous namespace + +void UnswizzleGobPermuteSSE2(std::span output, std::span input, u32 bytes_per_pixel, + u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, + u32 stride) { + switch (bytes_per_pixel) { + case 1: + return UnswizzleGobPermuteSSE2Impl<1>(output, input, width, height, depth, block_height, + block_depth, stride); + case 2: + return UnswizzleGobPermuteSSE2Impl<2>(output, input, width, height, depth, block_height, + block_depth, stride); + case 4: + return UnswizzleGobPermuteSSE2Impl<4>(output, input, width, height, depth, block_height, + block_depth, stride); + case 8: + return UnswizzleGobPermuteSSE2Impl<8>(output, input, width, height, depth, block_height, + block_depth, stride); + case 16: + return UnswizzleGobPermuteSSE2Impl<16>(output, input, width, height, depth, block_height, + block_depth, stride); + default: + UNREACHABLE(); + } +} + +} // namespace Tegra::Texture \ No newline at end of file diff --git a/src/video_core/textures/decoders_sse2.h b/src/video_core/textures/decoders_sse2.h new file mode 100644 index 0000000000..b52242bb64 --- /dev/null +++ b/src/video_core/textures/decoders_sse2.h @@ -0,0 +1,16 @@ +// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project +// SPDX-License-Identifier: GPL-3.0-or-later + +#pragma once + +#include + +#include "common/common_types.h" + +namespace Tegra::Texture { + +void UnswizzleGobPermuteSSE2(std::span output, std::span input, u32 bytes_per_pixel, + u32 width, u32 height, u32 depth, u32 block_height, u32 block_depth, + u32 stride); + +} // namespace Tegra::Texture