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