From a158960c20eb1534146821cc9e508656bb7ec92b Mon Sep 17 00:00:00 2001 From: shadowbeat070 <70843316+shadowbeat070@users.noreply.github.com> Date: Fri, 24 Jul 2026 19:13:02 +0200 Subject: [PATCH] feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal) (#592) * feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal) Move RDNA2 exact-XOR deswizzle (swizzle modes 5/9/24/27, 4bpp) off the CPU onto a GPU compute pass. GnmTiling.GetDetileParams resolves the shared addressing into DetileParams; the CPU fallback and both GPU kernels consume the same params so they never disagree. Vulkan (verified bit-exact on NVIDIA): SpirvFixedShaders.CreateDetileCompute hand-emits the SPIR-V kernel; VulkanDetilePass.RecordDetile records the dispatch into the async batch command buffer (never a blocking submit on the render thread) with transients retired via fence; VulkanDetileSelfTest (SHARPEMU_DETILE_SELFTEST=1) checks both entry points against the CPU detile. Metal (Mac-untested): detile_compute.msl (detile_cs) + MetalDetilePass mirror the Vulkan pass. The active Metal path CPU-detiles via the new GnmTiling.DetileWithParams when a texture arrives packaged (empty RgbaPixels + TiledSource/Detile), keeping Metal correct under default-on with no regression; wiring MetalDetilePass live is the remaining on-device step. Flags: GPU detile is default-on (SHARPEMU_GPU_DETILE=0 disables); [GPU-DETILE] diagnostics gated behind SHARPEMU_LOG_GPU_DETILE=1. Tests: 17 detile unit tests pass, incl. DetileWithParams and GetDetileParams each matching TryDetile bit-for-bit across all supported modes/bpp, plus a SPIR-V structural-validity test. * feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal) Move RDNA2 exact-XOR deswizzle (swizzle modes 5/9/24/27, 4bpp) off the CPU onto a GPU compute pass. GnmTiling.GetDetileParams resolves the shared addressing into DetileParams that the CPU fallback and both GPU kernels consume, so they never disagree; everything else keeps the CPU path. Vulkan (verified bit-exact on NVIDIA): SpirvFixedShaders.CreateDetileCompute hand-emits the kernel; VulkanDetilePass.RecordDetile records into the async batch command buffer (never a blocking submit on the render thread) with transients retired via fence, falling back to CPU detile on failure. VulkanDetileSelfTest (SHARPEMU_DETILE_SELFTEST=1) checks both entry points. Metal (Mac-untested): detile_compute.msl + MetalDetilePass mirror the Vulkan pass; the active Metal path CPU-detiles via GnmTiling.DetileWithParams so it stays correct under default-on. Wiring MetalDetilePass live is a follow-up. Flags: default-on (SHARPEMU_GPU_DETILE=0 disables); diagnostics behind SHARPEMU_LOG_GPU_DETILE=1. Adds 17 passing detile unit tests. * Fix: added support layered texture support for the GPU-Detiling. * Fix: Added support for BlockTable (1 / 4 / 8 (Morton/Z-order)) * feat: added support for 8 and 16 bpp (bytes per element) * Fixed a build failure specific to this branch --------- --- src/SharpEmu.Libs/Agc/AgcExports.cs | 160 ++++ src/SharpEmu.Libs/Agc/GnmTiling.cs | 184 +++++ src/SharpEmu.Libs/Gpu/GuestGpuTypes.cs | 10 +- .../Gpu/Metal/MetalDetilePass.cs | 313 ++++++++ .../Gpu/Metal/MetalVideoPresenter.Draws.cs | 40 + .../VideoOut/VulkanDetilePass.cs | 742 ++++++++++++++++++ .../VideoOut/VulkanDetileSelfTest.cs | 498 ++++++++++++ .../VideoOut/VulkanVideoPresenter.cs | 225 +++++- .../MslFixedShaders.cs | 8 + .../Templates/detile_compute.msl | 70 ++ .../SpirvFixedShaders.cs | 229 ++++++ .../Agc/DetileComputeSpirvTests.cs | 102 +++ .../Agc/GnmTilingDetileTests.cs | 148 ++++ 13 files changed, 2702 insertions(+), 27 deletions(-) create mode 100644 src/SharpEmu.Libs/Gpu/Metal/MetalDetilePass.cs create mode 100644 src/SharpEmu.Libs/VideoOut/VulkanDetilePass.cs create mode 100644 src/SharpEmu.Libs/VideoOut/VulkanDetileSelfTest.cs create mode 100644 src/SharpEmu.ShaderCompiler.Metal/Templates/detile_compute.msl create mode 100644 tests/SharpEmu.Libs.Tests/Agc/DetileComputeSpirvTests.cs diff --git a/src/SharpEmu.Libs/Agc/AgcExports.cs b/src/SharpEmu.Libs/Agc/AgcExports.cs index 0715c743..dfbb8cf6 100644 --- a/src/SharpEmu.Libs/Agc/AgcExports.cs +++ b/src/SharpEmu.Libs/Agc/AgcExports.cs @@ -271,6 +271,25 @@ public static partial class AgcExports Environment.GetEnvironmentVariable("SHARPEMU_NO_TEXTURE_SKIP"), "1", StringComparison.Ordinal); + + // GPU deswizzle: ship raw tiled bytes + params to the backend instead of + // detiling on the CPU. On by default; SHARPEMU_GPU_DETILE=0 forces the CPU + // path. Backend-agnostic here (only inspects DetileParams); the Vulkan/Metal + // backends detile on the GPU, others fall back to the CPU path. + private static readonly bool _gpuDetileEnabled = !string.Equals( + Environment.GetEnvironmentVariable("SHARPEMU_GPU_DETILE"), + "0", + StringComparison.Ordinal); + + // Diagnostics (SHARPEMU_LOG_GPU_DETILE=1): one line per distinct texture tile + // mode and per-gate decision, so we can see which swizzle modes/formats a + // title uses and whether each takes the GPU or CPU path. + private static readonly bool _gpuDetileLog = string.Equals( + Environment.GetEnvironmentVariable("SHARPEMU_LOG_GPU_DETILE"), + "1", + StringComparison.Ordinal); + private static readonly HashSet _seenTextureTileModes = new(); + private static readonly HashSet _gpuDetileGateDiag = new(); private static long _dcbWriteDataTraceCount; private static int _tracedVertexRangeCount; private static long _dcbWaitRegMemTraceCount; @@ -8235,6 +8254,15 @@ public static partial class AgcExports /// enabled and the format is understood; returns null to keep the raw /// bytes (linear surfaces, unknown modes, or non-power-of-two elements). /// + // The GPU detile kernel implements these two equation families at 4/8/16 bpp + // (one/two/four 32-bit words per element; 1/2 bpp are sub-word and stay on the + // CPU). Keep in lockstep with VulkanDetilePass.Supports / MetalDetilePass.Supports. + private static bool IsGpuDetileEquation(DetileEquation equation) => + equation == DetileEquation.ExactXor || equation == DetileEquation.BlockTable; + + private static bool IsGpuDetileBytesPerElement(int bytesPerElement) => + bytesPerElement is 4 or 8 or 16; + private static bool TryGetTextureElementLayout( TextureDescriptor descriptor, uint sourceWidth, @@ -8411,6 +8439,20 @@ public static partial class AgcExports return true; } + if (_gpuDetileLog) + { + lock (_seenTextureTileModes) + { + if (_seenTextureTileModes.Add(descriptor.TileMode)) + { + Console.Error.WriteLine( + $"[GPU-DETILE] texture tile_mode={descriptor.TileMode} fmt={descriptor.Format} " + + $"{descriptor.Width}x{descriptor.Height} " + + $"(0=linear; GPU covers exact-XOR 5/9/24/27 @ 4bpp)."); + } + } + } + var sourceWidth = descriptor.TileMode == 0 ? GetLinearTexturePitch( Math.Max(descriptor.Width, descriptor.Pitch), @@ -8695,6 +8737,65 @@ public static partial class AgcExports var arrayLayers = arrayUploadLayers; var layerBytes = checked((int)sourceSliceByteCount); var totalBytes = (long)layerBytes * arrayLayers; + + // GPU detile for arrayed exact-XOR/4bpp textures: pack the tiled array + // slices contiguously and hand them to the GPU pass (one dispatch-Z + // layer per slice), mirroring the single-layer gate above. The backend + // deswizzles every layer on the GPU; only unsupported cases fall to the + // CPU per-layer detile below. Font/text atlases uploaded as 2D arrays + // take this path. + if (_gpuDetileEnabled && hasElementLayout && !baseMipInTail && + IsGpuDetileBytesPerElement(bytesPerElement) && + (long)physicalSourceByteCount * arrayLayers <= int.MaxValue) + { + var gpuArrayParams = GnmTiling.GetDetileParams( + descriptor.TileMode, bytesPerElement, elementsWide, elementsHigh); + if (IsGpuDetileEquation(gpuArrayParams.Equation) && + (long)elementsWide * elementsHigh * bytesPerElement <= (long)physicalSourceByteCount) + { + var sliceBytes = checked((int)physicalSourceByteCount); + var tiledLayers = new byte[(long)sliceBytes * arrayLayers]; + var readAllLayers = true; + for (var layer = 0u; layer < arrayLayers; layer++) + { + if (!ctx.Memory.TryRead( + descriptor.Address + layer * chainSliceBytes + baseMipByteOffset, + tiledLayers.AsSpan(checked((int)(layer * (uint)sliceBytes)), sliceBytes))) + { + readAllLayers = false; + break; + } + } + + if (readAllLayers) + { + texture = new GuestDrawTexture( + descriptor.Address, + descriptor.Width, + descriptor.Height, + descriptor.Format, + descriptor.NumberType, + [], + IsFallback: false, + IsStorage: false, + MipLevels: descriptor.MipLevels, + MipLevel: mipLevel, + BaseMipLevel: descriptor.ViewBaseLevel, + ResourceMipLevels: descriptor.ResourceMipLevels, + Pitch: sourceWidth, + TileMode: descriptor.TileMode, + DstSelect: descriptor.DstSelect, + Sampler: sampler, + WriteGeneration: hasWriteGeneration ? writeGeneration : -1, + ArrayedView: true, + ArrayLayers: arrayLayers, + TiledSource: tiledLayers, + Detile: gpuArrayParams); + return true; + } + } + } + if (totalBytes <= int.MaxValue) { var layered = new byte[totalBytes]; @@ -8792,6 +8893,65 @@ public static partial class AgcExports } DumpTextureSourceIfRequested(descriptor, sourceWidth, source); + if (_gpuDetileLog && descriptor.TileMode != 0) + { + lock (_gpuDetileGateDiag) + { + if (_gpuDetileGateDiag.Add(descriptor.TileMode)) + { + var eq = hasElementLayout + ? GnmTiling.GetDetileParams( + descriptor.TileMode, bytesPerElement, elementsWide, elementsHigh).Equation + : DetileEquation.None; + Console.Error.WriteLine( + $"[GPU-DETILE] gate mode={descriptor.TileMode} fmt={descriptor.Format} " + + $"bpp={bytesPerElement} hasLayout={hasElementLayout} mipTail={baseMipInTail} " + + $"storage={isStorage} arrayed={isArrayed} eq={eq} -> " + + $"{(hasElementLayout && !baseMipInTail && IsGpuDetileBytesPerElement(bytesPerElement) && IsGpuDetileEquation(eq) ? "GPU" : "CPU")}"); + } + } + } + + // GPU detile: for the 4/8/16-bytes/element base-mip case the backend can + // deswizzle on the GPU (exact-XOR and block-table equations, including + // block-compressed formats), so ship the raw tiled bytes + params rather + // than paying the CPU detile. Everything else keeps the CPU path below. + // + // Arrayed textures are handled by the arrayed branch above (they package + // every layer's tiled slice); this branch is the single-layer case. + if (_gpuDetileEnabled && hasElementLayout && !baseMipInTail && + IsGpuDetileBytesPerElement(bytesPerElement) && !isArrayed) + { + var gpuDetileParams = GnmTiling.GetDetileParams( + descriptor.TileMode, bytesPerElement, elementsWide, elementsHigh); + if (IsGpuDetileEquation(gpuDetileParams.Equation) && + (long)elementsWide * elementsHigh * bytesPerElement <= source.Length) + { + texture = new GuestDrawTexture( + descriptor.Address, + descriptor.Width, + descriptor.Height, + descriptor.Format, + descriptor.NumberType, + [], + IsFallback: false, + IsStorage: isStorage, + MipLevels: descriptor.MipLevels, + MipLevel: mipLevel, + BaseMipLevel: descriptor.ViewBaseLevel, + ResourceMipLevels: descriptor.ResourceMipLevels, + Pitch: sourceWidth, + TileMode: descriptor.TileMode, + DstSelect: descriptor.DstSelect, + Sampler: ToGuestSampler(samplerDescriptor), + WriteGeneration: hasWriteGeneration ? writeGeneration : -1, + ArrayedView: isArrayed, + TiledSource: source, + Detile: gpuDetileParams); + return true; + } + } + var rgba = TryDetileTextureSource( descriptor, sourceWidth, diff --git a/src/SharpEmu.Libs/Agc/GnmTiling.cs b/src/SharpEmu.Libs/Agc/GnmTiling.cs index d0d1beef..515a1f75 100644 --- a/src/SharpEmu.Libs/Agc/GnmTiling.cs +++ b/src/SharpEmu.Libs/Agc/GnmTiling.cs @@ -6,6 +6,55 @@ using System.Runtime.CompilerServices; namespace SharpEmu.Libs.Agc; +/// Which in-block address equation a carries. +internal enum DetileEquation +{ + /// Unsupported mode/format; caller must use the CPU path or raw upload. + None, + + /// Exact AddrLib XOR equation (RDNA2 modes 5/9/24/27): factored X/Y terms. + ExactXor, + + /// Other modes: a precomputed in-block Morton/standard element-offset table. + BlockTable, +} + +/// +/// Backend-agnostic description of how to deswizzle one surface, produced by +/// . Holds only plain integers and small +/// int[] tables — no host graphics-API types — so it can cross the guest-GPU +/// backend seam and drive a Vulkan (SPIR-V) or Metal (MSL) detile compute kernel +/// identically to the CPU fallback. The single +/// shared addressing formula both consume is: +/// +/// inBlockByte = Equation == ExactXor +/// ? XByteTerm[x & XMask] ^ YByteTerm[y & YMask] +/// : BlockTable[(y % BlockHeight) * BlockWidth + (x % BlockWidth)] * BytesPerElement; +/// srcByte = ((y / BlockHeight) * BlocksPerRow + (x / BlockWidth)) * BlockBytes + inBlockByte; +/// +/// +internal readonly record struct DetileParams( + DetileEquation Equation, + int ElementsWide, + int ElementsHigh, + int BytesPerElement, + int BlockWidth, + int BlockHeight, + int BlockElements, + int BlockBytes, + int BlocksPerRow, + // ExactXor: within-block BYTE offset = XByteTerm[x & XMask] ^ YByteTerm[y & YMask]. + int[] XByteTerm, + int XMask, + int[] YByteTerm, + int YMask, + // BlockTable: within-block ELEMENT offset = BlockTable[inBlockY * BlockWidth + inBlockX]. + int[] BlockTable) +{ + /// False when the mode/format is not GPU-portable (Equation == None). + public bool IsSupported => Equation != DetileEquation.None; +} + /// /// Deswizzles RDNA2 (GFX10) tiled texture surfaces into linear layout so they /// can be uploaded to Vulkan. PS5 stores most textures in a swizzled layout @@ -501,6 +550,141 @@ internal static unsafe class GnmTiling return true; } + /// + /// Computes the detile parameters for a surface without performing the copy, + /// so a GPU compute kernel can run the deswizzle instead of the CPU. Returns + /// == false (Equation == None) when the + /// mode/format is not GPU-portable, so the caller keeps the CPU + /// path or a raw upload. Reuses the same helpers and + /// caches as , so the two never disagree on addressing. + /// + public static DetileParams GetDetileParams( + uint swizzleMode, + int bytesPerElement, + int elementsWide, + int elementsHigh) + { + if (!ShouldDetile(swizzleMode) || + bytesPerElement <= 0 || + elementsWide <= 0 || + elementsHigh <= 0 || + !TryGetSwizzleKind(swizzleMode, out var kind, out var blockBytes)) + { + return default; + } + + var bppLog2 = BitLog2((uint)bytesPerElement); + if (bppLog2 < 0) + { + return default; + } + + var blockElements = blockBytes >> bppLog2; + var (blockWidth, blockHeight) = SquareBlockDimensions(blockElements); + if (blockWidth == 0 || blockHeight == 0) + { + return default; + } + + var blocksPerRow = (elementsWide + blockWidth - 1) / blockWidth; + + if (TryGetExactXorPattern(swizzleMode, bppLog2, out var pattern)) + { + var terms = _patternTermCache.GetOrAdd( + (swizzleMode, bppLog2), + _ => CreatePatternTerms(pattern)); + return new DetileParams( + DetileEquation.ExactXor, + elementsWide, + elementsHigh, + bytesPerElement, + blockWidth, + blockHeight, + blockElements, + blockBytes, + blocksPerRow, + terms.X, + terms.XMask, + terms.Y, + terms.YMask, + []); + } + + var blockTable = _blockTableCache.GetOrAdd( + (kind, blockWidth, blockHeight), + static key => CreateBlockTable(key.Kind, key.Width, key.Height)); + return new DetileParams( + DetileEquation.BlockTable, + elementsWide, + elementsHigh, + bytesPerElement, + blockWidth, + blockHeight, + blockElements, + blockBytes, + blocksPerRow, + [], + 0, + [], + 0, + blockTable); + } + + /// + /// CPU deswizzle driven entirely by a resolved — the + /// exact addressing the Vulkan/Metal compute kernel runs per texel, so a + /// backend that packaged for the GPU path can + /// fall back to this without re-deriving the swizzle. Copies + /// ElementsWide * ElementsHigh elements from + /// into ; returns false when unsupported or the output + /// span is too small. Out-of-range source elements are left zero (matching the + /// reference), so a truncated degrades gracefully. + /// + public static bool DetileWithParams( + in DetileParams parameters, + ReadOnlySpan tiled, + Span linear) + { + if (!parameters.IsSupported) + { + return false; + } + + var width = parameters.ElementsWide; + var height = parameters.ElementsHigh; + var bpp = parameters.BytesPerElement; + var requiredLinear = (long)width * height * bpp; + if (width <= 0 || height <= 0 || bpp <= 0 || linear.Length < requiredLinear) + { + return false; + } + + var isExactXor = parameters.Equation == DetileEquation.ExactXor; + for (var y = 0; y < height; y++) + { + var blockY = y / parameters.BlockHeight; + var inY = y % parameters.BlockHeight; + var yTerm = isExactXor ? parameters.YByteTerm[y & parameters.YMask] : 0; + for (var x = 0; x < width; x++) + { + var blockX = x / parameters.BlockWidth; + var inBlockByte = isExactXor + ? parameters.XByteTerm[x & parameters.XMask] ^ yTerm + : parameters.BlockTable[inY * parameters.BlockWidth + (x % parameters.BlockWidth)] * bpp; + var srcByte = ((long)blockY * parameters.BlocksPerRow + blockX) * parameters.BlockBytes + inBlockByte; + var dstByte = ((long)y * width + x) * bpp; + if (srcByte < 0 || srcByte + bpp > tiled.Length) + { + continue; + } + + tiled.Slice((int)srcByte, bpp).CopyTo(linear.Slice((int)dstByte, bpp)); + } + } + + return true; + } + private enum SwizzleKind { Standard, diff --git a/src/SharpEmu.Libs/Gpu/GuestGpuTypes.cs b/src/SharpEmu.Libs/Gpu/GuestGpuTypes.cs index 6b4c7f86..5f883fe9 100644 --- a/src/SharpEmu.Libs/Gpu/GuestGpuTypes.cs +++ b/src/SharpEmu.Libs/Gpu/GuestGpuTypes.cs @@ -1,6 +1,8 @@ // Copyright (C) 2026 SharpEmu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +using SharpEmu.Libs.Agc; + namespace SharpEmu.Libs.Gpu; // The types that cross the guest-GPU backend seam. Every field is either a neutral @@ -35,7 +37,13 @@ internal sealed record GuestDrawTexture( bool ArrayedView = false, uint ArrayLayers = 1, uint Type = 9, - uint Depth = 1); + uint Depth = 1, + // GPU-detile opt-in (SHARPEMU_GPU_DETILE): when Detile is non-null the AGC + // layer skipped the CPU deswizzle and shipped the raw TILED bytes here in + // TiledSource; the Vulkan backend detiles them on the GPU. RgbaPixels is + // empty in that case. Both are neutral (no host graphics-API values). + byte[]? TiledSource = null, + DetileParams? Detile = null); /// Raw guest sampler descriptor dwords, copied verbatim from guest memory. internal readonly record struct GuestSampler( diff --git a/src/SharpEmu.Libs/Gpu/Metal/MetalDetilePass.cs b/src/SharpEmu.Libs/Gpu/Metal/MetalDetilePass.cs new file mode 100644 index 00000000..83da8d3d --- /dev/null +++ b/src/SharpEmu.Libs/Gpu/Metal/MetalDetilePass.cs @@ -0,0 +1,313 @@ +// Copyright (C) 2026 SharpEmu Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +using System.Numerics; +using SharpEmu.Libs.Agc; +using SharpEmu.ShaderCompiler.Metal; + +namespace SharpEmu.Libs.Gpu.Metal; + +/// +/// Metal twin of VulkanDetilePass: runs the ExactXor detile equation from +/// as a Metal compute kernel +/// (), writing a linear buffer +/// and blitting it into the sampled texture. +/// +/// records the compute dispatch + blit onto a caller's +/// command buffer and returns its transient buffers for the caller to release +/// once that command buffer completes — the async, non-blocking shape (Metal +/// hazard-tracks the compute-write → blit-read → sample dependency automatically, +/// so no manual barriers are needed). +/// +/// Only ExactXor 4-bytes/element surfaces are handled. NOTE: authored on Windows; +/// the MSL and every Metal call here are Mac-untested — mirrors the +/// verified Vulkan logic and the existing Metal message-send conventions, but +/// must be validated on a real Metal device. +/// +internal sealed unsafe class MetalDetilePass : IDisposable +{ + private const uint LocalSize = 8; + private const int PushConstantUints = 11; + + private readonly nint _device; + private nint _pipelineState; + private bool _initialized; + private bool _disposed; + + public MetalDetilePass(nint device) + { + _device = device; + } + + public static bool Supports(in DetileParams parameters) => + (parameters.Equation == DetileEquation.ExactXor || + parameters.Equation == DetileEquation.BlockTable) && + parameters.BytesPerElement is 4 or 8 or 16; + + /// + /// Records the deswizzle of into + /// ( x + /// texels x slices) + /// onto . The kernel iterates the element grid + /// from (for block-compressed formats a 4x4 block + /// is one element). Does not commit; the caller releases + /// when the command buffer completes. + /// Returns false (empty transients) when unsupported or the pipeline could not + /// be built. + /// + public bool RecordDetile( + nint commandBuffer, + nint texture, + uint texelWidth, + uint texelHeight, + uint layers, + ReadOnlySpan tiled, + in DetileParams parameters, + out nint[] transientBuffers) + { + transientBuffers = []; + var bytesPerElement = (uint)parameters.BytesPerElement; + if (_disposed || commandBuffer == 0 || texture == 0 || + !Supports(parameters) || texelWidth == 0 || texelHeight == 0 || layers == 0 || tiled.IsEmpty || + tiled.Length % (int)(layers * bytesPerElement) != 0 || + !EnsurePipeline()) + { + return false; + } + + var elementsWide = (uint)parameters.ElementsWide; + var elementsHigh = (uint)parameters.ElementsHigh; + var uintsPerElement = bytesPerElement / sizeof(uint); + + // Array slices are packed contiguously in the tiled buffer; each slice's + // element stride is the whole buffer split evenly by layer. + var srcSliceElements = (uint)((ulong)tiled.Length / bytesPerElement / layers); + + // Binding 1 carries the within-block offset table. ExactXor: element-shifted + // X/Y byte terms. BlockTable: GetDetileParams' block table (already element + // offsets) in binding 1, a placeholder in binding 2. The two equations index + // different-sized buffers, so the kernel branches and reads only one. + uint[] xTerm; + uint[] yTerm; + uint equationValue; + if (parameters.Equation == DetileEquation.BlockTable) + { + xTerm = new uint[parameters.BlockTable.Length]; + for (var index = 0; index < xTerm.Length; index++) + { + xTerm[index] = (uint)parameters.BlockTable[index]; + } + + yTerm = [0]; + equationValue = 1; + } + else + { + var shift = BitOperations.TrailingZeroCount((uint)parameters.BytesPerElement); + xTerm = ToElementTerms(parameters.XByteTerm, shift); + yTerm = ToElementTerms(parameters.YByteTerm, shift); + equationValue = 0; + } + + var newBufferWithBytes = MetalNative.Selector("newBufferWithBytes:length:options:"); + var newBufferWithLength = MetalNative.Selector("newBufferWithLength:options:"); + + nint tiledBuffer; + nint xBuffer; + nint yBuffer; + fixed (byte* tiledPointer = tiled) + { + tiledBuffer = MetalNative.SendBuffer( + _device, newBufferWithBytes, (nint)tiledPointer, (nuint)tiled.Length, 0); + } + + fixed (uint* xPointer = xTerm) + { + xBuffer = MetalNative.SendBuffer( + _device, newBufferWithBytes, (nint)xPointer, (nuint)xTerm.Length * sizeof(uint), 0); + } + + fixed (uint* yPointer = yTerm) + { + yBuffer = MetalNative.SendBuffer( + _device, newBufferWithBytes, (nint)yPointer, (nuint)yTerm.Length * sizeof(uint), 0); + } + + var outputBytes = (nuint)elementsWide * elementsHigh * bytesPerElement * layers; + var outputBuffer = MetalNative.SendNewBuffer(_device, newBufferWithLength, outputBytes, 0); + + Span push = + [ + elementsWide, + elementsHigh, + (uint)parameters.BlockWidth, + (uint)parameters.BlockHeight, + (uint)parameters.BlockElements, + (uint)parameters.BlocksPerRow, + (uint)parameters.XMask, + (uint)parameters.YMask, + srcSliceElements, + equationValue, + uintsPerElement, + ]; + nint paramsBuffer; + fixed (uint* pushPointer = push) + { + paramsBuffer = MetalNative.SendBuffer( + _device, newBufferWithBytes, (nint)pushPointer, (nuint)PushConstantUints * sizeof(uint), 0); + } + + if (tiledBuffer == 0 || xBuffer == 0 || yBuffer == 0 || outputBuffer == 0 || paramsBuffer == 0) + { + ReleaseAll(tiledBuffer, xBuffer, yBuffer, outputBuffer, paramsBuffer); + return false; + } + + // Compute encoder: one thread per texel. + var setBuffer = MetalNative.Selector("setBuffer:offset:atIndex:"); + var encoder = MetalNative.Send(commandBuffer, MetalNative.Selector("computeCommandEncoder")); + MetalNative.Send(encoder, MetalNative.Selector("setComputePipelineState:"), _pipelineState); + MetalNative.SendSetBuffer(encoder, setBuffer, tiledBuffer, 0, 0); + MetalNative.SendSetBuffer(encoder, setBuffer, xBuffer, 0, 1); + MetalNative.SendSetBuffer(encoder, setBuffer, yBuffer, 0, 2); + MetalNative.SendSetBuffer(encoder, setBuffer, outputBuffer, 0, 3); + MetalNative.SendSetBuffer(encoder, setBuffer, paramsBuffer, 0, 4); + + // X is widened by uintsPerElement (each thread copies one word); one + // grid-Z layer per array slice. + var threadgroups = new MtlSize + { + Width = (nuint)((elementsWide * uintsPerElement + LocalSize - 1) / LocalSize), + Height = (nuint)((elementsHigh + LocalSize - 1) / LocalSize), + Depth = layers, + }; + var threadsPerThreadgroup = new MtlSize { Width = LocalSize, Height = LocalSize, Depth = 1 }; + MetalNative.SendDispatch( + encoder, + MetalNative.Selector("dispatchThreadgroups:threadsPerThreadgroup:"), + threadgroups, + threadsPerThreadgroup); + MetalNative.SendVoid(encoder, MetalNative.Selector("endEncoding")); + + // Blit the layer-major linear output buffer into the sampled texture, one + // slice per array layer (Metal copyFromBuffer targets a single slice). The + // buffer is element/block-packed (row stride = elementsWide*bpp); the copy + // region is in texels. Metal tracks the compute-write -> blit-read hazard. + var blit = MetalNative.Send(commandBuffer, MetalNative.Selector("blitCommandEncoder")); + var copySelector = MetalNative.Selector( + "copyFromBuffer:sourceOffset:sourceBytesPerRow:sourceBytesPerImage:sourceSize:" + + "toTexture:destinationSlice:destinationLevel:destinationOrigin:"); + var sliceBytes = (nuint)elementsWide * elementsHigh * bytesPerElement; + var rowBytes = (nuint)elementsWide * bytesPerElement; + for (uint layer = 0; layer < layers; layer++) + { + MetalNative.SendCopyBufferToTexture( + blit, + copySelector, + outputBuffer, + (nuint)layer * sliceBytes, + rowBytes, + sliceBytes, + new MtlSize { Width = texelWidth, Height = texelHeight, Depth = 1 }, + texture, + layer, + 0, + new MtlOrigin { X = 0, Y = 0, Z = 0 }); + } + + MetalNative.SendVoid(blit, MetalNative.Selector("endEncoding")); + + transientBuffers = [tiledBuffer, xBuffer, yBuffer, outputBuffer, paramsBuffer]; + return true; + } + + private bool EnsurePipeline() + { + if (_initialized) + { + return _pipelineState != 0; + } + + _initialized = true; + + var options = MetalNative.Send( + MetalNative.Send(MetalNative.Class("MTLCompileOptions"), MetalNative.Selector("alloc")), + MetalNative.Selector("init")); + MetalNative.SendVoidBool(options, MetalNative.Selector("setFastMathEnabled:"), false); + + nint libraryError = 0; + var library = MetalNative.Send( + _device, + MetalNative.Selector("newLibraryWithSource:options:error:"), + MetalNative.NsString(MslFixedShaders.CreateDetileCompute()), + options, + ref libraryError); + if (library == 0) + { + Console.Error.WriteLine( + $"[GPU-DETILE] Metal detile library compile failed: {MetalNative.DescribeError(libraryError)}"); + return false; + } + + var function = MetalNative.Send( + library, MetalNative.Selector("newFunctionWithName:"), MetalNative.NsString("detile_cs")); + if (function == 0) + { + Console.Error.WriteLine("[GPU-DETILE] Metal detile function 'detile_cs' not found."); + return false; + } + + nint pipelineError = 0; + _pipelineState = MetalNative.Send( + _device, + MetalNative.Selector("newComputePipelineStateWithFunction:error:"), + function, + ref pipelineError); + if (_pipelineState == 0) + { + Console.Error.WriteLine( + $"[GPU-DETILE] Metal detile pipeline failed: {MetalNative.DescribeError(pipelineError)}"); + return false; + } + + return true; + } + + private static uint[] ToElementTerms(int[] byteTerms, int shift) + { + var terms = new uint[byteTerms.Length]; + for (var index = 0; index < byteTerms.Length; index++) + { + terms[index] = (uint)byteTerms[index] >> shift; + } + + return terms; + } + + private static void ReleaseAll(params nint[] objects) + { + var release = MetalNative.Selector("release"); + foreach (var handle in objects) + { + if (handle != 0) + { + MetalNative.SendVoid(handle, release); + } + } + } + + public void Dispose() + { + if (_disposed) + { + return; + } + + _disposed = true; + if (_pipelineState != 0) + { + MetalNative.SendVoid(_pipelineState, MetalNative.Selector("release")); + _pipelineState = 0; + } + } +} diff --git a/src/SharpEmu.Libs/Gpu/Metal/MetalVideoPresenter.Draws.cs b/src/SharpEmu.Libs/Gpu/Metal/MetalVideoPresenter.Draws.cs index 86740ef0..788e6227 100644 --- a/src/SharpEmu.Libs/Gpu/Metal/MetalVideoPresenter.Draws.cs +++ b/src/SharpEmu.Libs/Gpu/Metal/MetalVideoPresenter.Draws.cs @@ -1,6 +1,7 @@ // Copyright (C) 2026 SharpEmu Emulator Project // SPDX-License-Identifier: GPL-2.0-or-later +using SharpEmu.Libs.Agc; using SharpEmu.ShaderCompiler; using SharpEmu.ShaderCompiler.Metal; @@ -1583,6 +1584,45 @@ internal static partial class MetalVideoPresenter return cached; } + // Default-on GPU detile packages the tiled source + resolved DetileParams + // with empty RgbaPixels so a backend can deswizzle on the GPU. The Metal + // GPU compute pass (MetalDetilePass / detile_compute.msl) is the intended + // equivalent of VulkanDetilePass, but it is Mac-untested, so the active + // Metal path CPU-detiles here via GnmTiling.DetileWithParams — the exact + // same DetileParams addressing the kernel runs — restoring the linear-upload + // behavior Metal had before GPU detile existed, with no regression. + if (texture.RgbaPixels.Length == 0 && + texture.TiledSource is { } tiledSource && + texture.Detile is { } detileParameters) + { + // The tiled source packs the array slices contiguously (one per layer); + // detile each into its layer-major linear region so the reconstructed + // pixels match what the CPU array-upload path produced pre-GPU-detile. + // A plain 2D texture is just one layer. + var layers = Math.Max((int)texture.ArrayLayers, 1); + var sliceLinearBytes = + detileParameters.ElementsWide * detileParameters.ElementsHigh * detileParameters.BytesPerElement; + var sliceTiledBytes = tiledSource.Length / layers; + var linear = new byte[sliceLinearBytes * layers]; + var detiledAll = true; + for (var layer = 0; layer < layers; layer++) + { + if (!GnmTiling.DetileWithParams( + detileParameters, + tiledSource.AsSpan(layer * sliceTiledBytes, sliceTiledBytes), + linear.AsSpan(layer * sliceLinearBytes, sliceLinearBytes))) + { + detiledAll = false; + break; + } + } + + if (detiledAll) + { + texture = texture with { RgbaPixels = linear }; + } + } + // AGC ships the raw (detiled) source texels; create the texture in the // guest's native format — Mac-family GPUs sample BC blocks directly — // and size expectations with the same block-aware math AGC used. diff --git a/src/SharpEmu.Libs/VideoOut/VulkanDetilePass.cs b/src/SharpEmu.Libs/VideoOut/VulkanDetilePass.cs new file mode 100644 index 00000000..d3773dec --- /dev/null +++ b/src/SharpEmu.Libs/VideoOut/VulkanDetilePass.cs @@ -0,0 +1,742 @@ +// Copyright (C) 2026 SharpEmu Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +using System.Numerics; +using SharpEmu.Libs.Agc; +using SharpEmu.ShaderCompiler.Vulkan; +using Silk.NET.Vulkan; +using VkBuffer = Silk.NET.Vulkan.Buffer; + +namespace SharpEmu.Libs.VideoOut; + +/// +/// Self-contained GPU deswizzle pass: runs the ExactXor detile equation from +/// as a Vulkan compute shader +/// (), writing a linear buffer +/// and copying it into a sampled image — the GPU equivalent of the CPU +/// GnmTiling.TryDetile + staging upload. +/// +/// Two entry points share the same (verified) recording: +/// is a self-contained one-shot (submit + wait) used +/// by the isolation self-test; records into a caller's +/// command buffer and hands back its transient buffers + descriptor pool for the +/// caller to retire with that command buffer's fence — the render-path variant, +/// which must never block the render thread. +/// +/// Only ExactXor 4-bytes/element surfaces are handled; lets +/// the caller fall back to the CPU path for everything else. +/// +internal sealed unsafe class VulkanDetilePass : IDisposable +{ + private const uint LocalSize = 8; + private const uint PushConstantBytes = 11 * sizeof(uint); + + private readonly Vk _vk; + private readonly Device _device; + private readonly Queue _queue; + private readonly PhysicalDevice _physicalDevice; + private readonly uint _queueFamilyIndex; + + private ShaderModule _shaderModule; + private DescriptorSetLayout _descriptorSetLayout; + private PipelineLayout _pipelineLayout; + private Pipeline _pipeline; + private CommandPool _commandPool; + private bool _initialized; + private bool _disposed; + + public VulkanDetilePass( + Vk vk, + Device device, + Queue queue, + PhysicalDevice physicalDevice, + uint queueFamilyIndex) + { + _vk = vk; + _device = device; + _queue = queue; + _physicalDevice = physicalDevice; + _queueFamilyIndex = queueFamilyIndex; + } + + /// + /// The kernel handles the exact-XOR and block-table modes at 4/8/16 + /// bytes-per-element (one, two, or four 32-bit words per element). 1/2 bpp are + /// sub-word and stay on the CPU. + /// + public static bool Supports(in DetileParams parameters) => + (parameters.Equation == DetileEquation.ExactXor || + parameters.Equation == DetileEquation.BlockTable) && + parameters.BytesPerElement is 4 or 8 or 16; + + /// Transient per-detile resources the caller must retire once the + /// command buffer they were recorded into has completed. + public readonly record struct Transients( + (VkBuffer Buffer, DeviceMemory Memory)[] Buffers, + DescriptorPool DescriptorPool); + + private struct DetileResources + { + public VkBuffer Tiled; + public DeviceMemory TiledMemory; + public VkBuffer XTerm; + public DeviceMemory XMemory; + public VkBuffer YTerm; + public DeviceMemory YMemory; + public VkBuffer Output; + public DeviceMemory OutputMemory; + public DescriptorPool Pool; + public DescriptorSet Set; + public ulong OutputBytes; + public uint SrcSliceElements; + public uint EquationValue; + public uint UintsPerElement; + } + + /// + /// Records the deswizzle of into + /// ( x texels x + /// array slices, currently in + /// ) onto , + /// leaving the image . The kernel + /// iterates the element grid from (for + /// block-compressed formats a 4x4 block is one element, so the element grid is + /// smaller than the texel grid). The tiled buffer holds the array slices packed + /// contiguously (each an independently tiled 2D surface). Does not submit; the + /// caller retires with the command buffer's fence. + /// Returns false (with empty transients) when unsupported. + /// + public bool RecordDetile( + CommandBuffer commandBuffer, + Image image, + ImageLayout currentLayout, + uint texelWidth, + uint texelHeight, + uint layers, + ReadOnlySpan tiled, + in DetileParams parameters, + out Transients transients) + { + transients = new Transients([], default); + if (_disposed || !Supports(parameters) || texelWidth == 0 || texelHeight == 0 || layers == 0 || + tiled.IsEmpty || tiled.Length % (int)(layers * (uint)parameters.BytesPerElement) != 0) + { + return false; + } + + EnsurePipeline(); + + var resources = default(DetileResources); + try + { + PrepareResources(tiled, parameters, layers, ref resources); + RecordCommands(commandBuffer, in resources, image, currentLayout, texelWidth, texelHeight, layers, in parameters); + } + catch + { + DestroyResources(in resources); + throw; + } + + transients = new Transients( + [ + (resources.Tiled, resources.TiledMemory), + (resources.XTerm, resources.XMemory), + (resources.YTerm, resources.YMemory), + (resources.Output, resources.OutputMemory), + ], + resources.Pool); + return true; + } + + /// + /// One-shot variant used by the isolation self-test: records the detile onto a + /// private command buffer, submits, waits, and frees every transient. Never + /// call this on the render thread — its blocking wait would deadlock the + /// present pipeline; use there. + /// + public bool DetileIntoImage( + Image image, + ImageLayout currentLayout, + uint texelWidth, + uint texelHeight, + uint layers, + ReadOnlySpan tiled, + in DetileParams parameters) + { + if (_disposed || !Supports(parameters) || texelWidth == 0 || texelHeight == 0 || layers == 0 || + tiled.IsEmpty || tiled.Length % (int)(layers * (uint)parameters.BytesPerElement) != 0) + { + return false; + } + + EnsurePipeline(); + + var resources = default(DetileResources); + CommandBuffer commandBuffer = default; + Fence fence = default; + try + { + PrepareResources(tiled, parameters, layers, ref resources); + + commandBuffer = AllocateCommandBuffer(); + BeginCommandBuffer(commandBuffer); + RecordCommands(commandBuffer, in resources, image, currentLayout, texelWidth, texelHeight, layers, in parameters); + Check(_vk.EndCommandBuffer(commandBuffer), "vkEndCommandBuffer(detile)"); + + fence = CreateFence(); + var submitInfo = new SubmitInfo + { + SType = StructureType.SubmitInfo, + CommandBufferCount = 1, + PCommandBuffers = &commandBuffer, + }; + Check(_vk.QueueSubmit(_queue, 1, &submitInfo, fence), "vkQueueSubmit(detile)"); + Check(_vk.WaitForFences(_device, 1, &fence, true, ulong.MaxValue), "vkWaitForFences(detile)"); + return true; + } + finally + { + if (fence.Handle != 0) + { + _vk.DestroyFence(_device, fence, null); + } + + if (commandBuffer.Handle != 0) + { + _vk.FreeCommandBuffers(_device, _commandPool, 1, &commandBuffer); + } + + DestroyResources(in resources); + } + } + + private void PrepareResources(ReadOnlySpan tiled, in DetileParams parameters, uint layers, ref DetileResources resources) + { + // Binding 1 carries the within-block offset table, binding 2 the Y terms. + // ExactXor: xTerm/yTerm are byte offsets; the kernel indexes a uint[], so it + // wants element offsets — for a power-of-two element size the low + // log2(bpp) bits of every term are 0, so the right shift is exact. + // BlockTable: GetDetileParams' block table is already element offsets; it + // goes in binding 1 and binding 2 is an unused placeholder. + uint[] xTerm; + uint[] yTerm; + if (parameters.Equation == DetileEquation.BlockTable) + { + xTerm = new uint[parameters.BlockTable.Length]; + for (var index = 0; index < xTerm.Length; index++) + { + xTerm[index] = (uint)parameters.BlockTable[index]; + } + + yTerm = [0]; + resources.EquationValue = 1; + } + else + { + var shift = BitOperations.TrailingZeroCount((uint)parameters.BytesPerElement); + xTerm = ToElementTerms(parameters.XByteTerm, shift); + yTerm = ToElementTerms(parameters.YByteTerm, shift); + resources.EquationValue = 0; + } + + // The array slices are packed contiguously in the tiled buffer, so each + // slice's element stride is the whole tiled buffer split evenly by layer. + // Element sizes are in bytes-per-element; the kernel moves bpp/4 words each. + var bytesPerElement = (uint)parameters.BytesPerElement; + resources.UintsPerElement = bytesPerElement / sizeof(uint); + resources.SrcSliceElements = (uint)((ulong)tiled.Length / bytesPerElement / layers); + resources.OutputBytes = + (ulong)parameters.ElementsWide * (ulong)parameters.ElementsHigh * bytesPerElement * layers; + + resources.Tiled = CreateHostBuffer((ulong)tiled.Length, BufferUsageFlags.StorageBufferBit, out resources.TiledMemory); + UploadBytes(resources.TiledMemory, tiled); + resources.XTerm = CreateHostBuffer((ulong)xTerm.Length * sizeof(uint), BufferUsageFlags.StorageBufferBit, out resources.XMemory); + UploadUInts(resources.XMemory, xTerm); + resources.YTerm = CreateHostBuffer((ulong)yTerm.Length * sizeof(uint), BufferUsageFlags.StorageBufferBit, out resources.YMemory); + UploadUInts(resources.YMemory, yTerm); + resources.Output = CreateHostBuffer( + resources.OutputBytes, + BufferUsageFlags.StorageBufferBit | BufferUsageFlags.TransferSrcBit, + out resources.OutputMemory); + + resources.Pool = CreateDescriptorPool(); + resources.Set = AllocateDescriptorSet(resources.Pool); + WriteDescriptors( + resources.Set, + (resources.Tiled, (ulong)tiled.Length), + (resources.XTerm, (ulong)xTerm.Length * sizeof(uint)), + (resources.YTerm, (ulong)yTerm.Length * sizeof(uint)), + (resources.Output, resources.OutputBytes)); + } + + private void RecordCommands( + CommandBuffer commandBuffer, + in DetileResources resources, + Image image, + ImageLayout currentLayout, + uint texelWidth, + uint texelHeight, + uint layers, + in DetileParams parameters) + { + // The kernel iterates the element grid (smaller than the texel grid for + // block-compressed formats); the image copy below uses the texel grid. + var elementsWide = (uint)parameters.ElementsWide; + var elementsHigh = (uint)parameters.ElementsHigh; + + var descriptorSet = resources.Set; + _vk.CmdBindPipeline(commandBuffer, PipelineBindPoint.Compute, _pipeline); + _vk.CmdBindDescriptorSets( + commandBuffer, PipelineBindPoint.Compute, _pipelineLayout, 0, 1, &descriptorSet, 0, null); + + Span push = + [ + elementsWide, + elementsHigh, + (uint)parameters.BlockWidth, + (uint)parameters.BlockHeight, + (uint)parameters.BlockElements, + (uint)parameters.BlocksPerRow, + (uint)parameters.XMask, + (uint)parameters.YMask, + resources.SrcSliceElements, + resources.EquationValue, + resources.UintsPerElement, + ]; + fixed (uint* pushPointer = push) + { + _vk.CmdPushConstants( + commandBuffer, _pipelineLayout, ShaderStageFlags.ComputeBit, 0, PushConstantBytes, pushPointer); + } + + // X is widened by uintsPerElement (each thread copies one word); one + // dispatch-Z layer per array slice. + _vk.CmdDispatch( + commandBuffer, + (elementsWide * resources.UintsPerElement + LocalSize - 1) / LocalSize, + (elementsHigh + LocalSize - 1) / LocalSize, + layers); + + // Compute store -> transfer read on the linear output buffer. + var outputBarrier = new BufferMemoryBarrier + { + SType = StructureType.BufferMemoryBarrier, + SrcAccessMask = AccessFlags.ShaderWriteBit, + DstAccessMask = AccessFlags.TransferReadBit, + SrcQueueFamilyIndex = Vk.QueueFamilyIgnored, + DstQueueFamilyIndex = Vk.QueueFamilyIgnored, + Buffer = resources.Output, + Offset = 0, + Size = resources.OutputBytes, + }; + _vk.CmdPipelineBarrier( + commandBuffer, + PipelineStageFlags.ComputeShaderBit, + PipelineStageFlags.TransferBit, + 0, + 0, + null, + 1, + &outputBarrier, + 0, + null); + + var initialized = currentLayout == ImageLayout.ShaderReadOnlyOptimal; + TransitionImage( + commandBuffer, + image, + currentLayout, + ImageLayout.TransferDstOptimal, + initialized ? AccessFlags.ShaderReadBit : 0, + AccessFlags.TransferWriteBit, + initialized ? PipelineStageFlags.FragmentShaderBit : PipelineStageFlags.TopOfPipeBit, + PipelineStageFlags.TransferBit, + layers); + + // The output buffer is layer-major, tightly packed (BufferRowLength 0 => + // one element-row per texel-row, which for compressed formats is the block + // row), so a single copy fills every array layer. Extent is in texels. + var copyRegion = new BufferImageCopy + { + BufferOffset = 0, + BufferRowLength = 0, + BufferImageHeight = 0, + ImageSubresource = new ImageSubresourceLayers(ImageAspectFlags.ColorBit, 0, 0, layers), + ImageOffset = default, + ImageExtent = new Extent3D(texelWidth, texelHeight, 1), + }; + _vk.CmdCopyBufferToImage( + commandBuffer, resources.Output, image, ImageLayout.TransferDstOptimal, 1, ©Region); + + TransitionImage( + commandBuffer, + image, + ImageLayout.TransferDstOptimal, + ImageLayout.ShaderReadOnlyOptimal, + AccessFlags.TransferWriteBit, + AccessFlags.ShaderReadBit, + PipelineStageFlags.TransferBit, + PipelineStageFlags.FragmentShaderBit, + layers); + } + + private void DestroyResources(in DetileResources resources) + { + if (resources.Pool.Handle != 0) + { + _vk.DestroyDescriptorPool(_device, resources.Pool, null); + } + + DestroyBuffer(resources.Output, resources.OutputMemory); + DestroyBuffer(resources.YTerm, resources.YMemory); + DestroyBuffer(resources.XTerm, resources.XMemory); + DestroyBuffer(resources.Tiled, resources.TiledMemory); + } + + private void EnsurePipeline() + { + if (_initialized) + { + return; + } + + var spirv = SpirvFixedShaders.CreateDetileCompute(); + fixed (byte* code = spirv) + { + var moduleInfo = new ShaderModuleCreateInfo + { + SType = StructureType.ShaderModuleCreateInfo, + CodeSize = (nuint)spirv.Length, + PCode = (uint*)code, + }; + Check( + _vk.CreateShaderModule(_device, &moduleInfo, null, out _shaderModule), + "vkCreateShaderModule(detile)"); + } + + var bindings = stackalloc DescriptorSetLayoutBinding[4]; + for (uint index = 0; index < 4; index++) + { + bindings[index] = new DescriptorSetLayoutBinding + { + Binding = index, + DescriptorType = DescriptorType.StorageBuffer, + DescriptorCount = 1, + StageFlags = ShaderStageFlags.ComputeBit, + }; + } + + var layoutInfo = new DescriptorSetLayoutCreateInfo + { + SType = StructureType.DescriptorSetLayoutCreateInfo, + BindingCount = 4, + PBindings = bindings, + }; + Check( + _vk.CreateDescriptorSetLayout(_device, &layoutInfo, null, out _descriptorSetLayout), + "vkCreateDescriptorSetLayout(detile)"); + + var pushRange = new PushConstantRange + { + StageFlags = ShaderStageFlags.ComputeBit, + Offset = 0, + Size = PushConstantBytes, + }; + var setLayout = _descriptorSetLayout; + var pipelineLayoutInfo = new PipelineLayoutCreateInfo + { + SType = StructureType.PipelineLayoutCreateInfo, + SetLayoutCount = 1, + PSetLayouts = &setLayout, + PushConstantRangeCount = 1, + PPushConstantRanges = &pushRange, + }; + Check( + _vk.CreatePipelineLayout(_device, &pipelineLayoutInfo, null, out _pipelineLayout), + "vkCreatePipelineLayout(detile)"); + + ReadOnlySpan entryPoint = "main\0"u8; + fixed (byte* entry = entryPoint) + { + var pipelineInfo = new ComputePipelineCreateInfo + { + SType = StructureType.ComputePipelineCreateInfo, + Layout = _pipelineLayout, + Stage = new PipelineShaderStageCreateInfo + { + SType = StructureType.PipelineShaderStageCreateInfo, + Stage = ShaderStageFlags.ComputeBit, + Module = _shaderModule, + PName = entry, + }, + }; + Check( + _vk.CreateComputePipelines(_device, default, 1, &pipelineInfo, null, out _pipeline), + "vkCreateComputePipelines(detile)"); + } + + var poolInfo = new CommandPoolCreateInfo + { + SType = StructureType.CommandPoolCreateInfo, + QueueFamilyIndex = _queueFamilyIndex, + Flags = CommandPoolCreateFlags.ResetCommandBufferBit, + }; + Check( + _vk.CreateCommandPool(_device, &poolInfo, null, out _commandPool), + "vkCreateCommandPool(detile)"); + + _initialized = true; + } + + private static uint[] ToElementTerms(int[] byteTerms, int shift) + { + var terms = new uint[byteTerms.Length]; + for (var index = 0; index < byteTerms.Length; index++) + { + terms[index] = (uint)byteTerms[index] >> shift; + } + + return terms; + } + + private VkBuffer CreateHostBuffer(ulong size, BufferUsageFlags usage, out DeviceMemory memory) + { + var bufferInfo = new BufferCreateInfo + { + SType = StructureType.BufferCreateInfo, + Size = size, + Usage = usage, + SharingMode = SharingMode.Exclusive, + }; + Check(_vk.CreateBuffer(_device, &bufferInfo, null, out var buffer), "vkCreateBuffer(detile)"); + + _vk.GetBufferMemoryRequirements(_device, buffer, out var requirements); + var allocateInfo = new MemoryAllocateInfo + { + SType = StructureType.MemoryAllocateInfo, + AllocationSize = requirements.Size, + MemoryTypeIndex = FindMemoryType( + requirements.MemoryTypeBits, + MemoryPropertyFlags.HostVisibleBit | MemoryPropertyFlags.HostCoherentBit), + }; + Check(_vk.AllocateMemory(_device, &allocateInfo, null, out memory), "vkAllocateMemory(detile)"); + Check(_vk.BindBufferMemory(_device, buffer, memory, 0), "vkBindBufferMemory(detile)"); + return buffer; + } + + private uint FindMemoryType(uint typeBits, MemoryPropertyFlags requiredFlags) + { + _vk.GetPhysicalDeviceMemoryProperties(_physicalDevice, out var properties); + var memoryTypes = &properties.MemoryTypes.Element0; + for (uint index = 0; index < properties.MemoryTypeCount; index++) + { + if ((typeBits & (1u << (int)index)) != 0 && + (memoryTypes[index].PropertyFlags & requiredFlags) == requiredFlags) + { + return index; + } + } + + throw new InvalidOperationException("No compatible Vulkan host-visible memory type for detile."); + } + + private void UploadBytes(DeviceMemory memory, ReadOnlySpan data) + { + void* mapped; + Check(_vk.MapMemory(_device, memory, 0, (ulong)data.Length, 0, &mapped), "vkMapMemory(detile)"); + data.CopyTo(new Span(mapped, data.Length)); + _vk.UnmapMemory(_device, memory); + } + + private void UploadUInts(DeviceMemory memory, uint[] data) + { + void* mapped; + var byteCount = (ulong)data.Length * sizeof(uint); + Check(_vk.MapMemory(_device, memory, 0, byteCount, 0, &mapped), "vkMapMemory(detile terms)"); + data.AsSpan().CopyTo(new Span(mapped, data.Length)); + _vk.UnmapMemory(_device, memory); + } + + private DescriptorPool CreateDescriptorPool() + { + var poolSize = new DescriptorPoolSize + { + Type = DescriptorType.StorageBuffer, + DescriptorCount = 4, + }; + var poolInfo = new DescriptorPoolCreateInfo + { + SType = StructureType.DescriptorPoolCreateInfo, + MaxSets = 1, + PoolSizeCount = 1, + PPoolSizes = &poolSize, + }; + Check( + _vk.CreateDescriptorPool(_device, &poolInfo, null, out var pool), + "vkCreateDescriptorPool(detile)"); + return pool; + } + + private DescriptorSet AllocateDescriptorSet(DescriptorPool pool) + { + var setLayout = _descriptorSetLayout; + var allocateInfo = new DescriptorSetAllocateInfo + { + SType = StructureType.DescriptorSetAllocateInfo, + DescriptorPool = pool, + DescriptorSetCount = 1, + PSetLayouts = &setLayout, + }; + Check( + _vk.AllocateDescriptorSets(_device, &allocateInfo, out var descriptorSet), + "vkAllocateDescriptorSets(detile)"); + return descriptorSet; + } + + private void WriteDescriptors( + DescriptorSet descriptorSet, + (VkBuffer Buffer, ulong Size) binding0, + (VkBuffer Buffer, ulong Size) binding1, + (VkBuffer Buffer, ulong Size) binding2, + (VkBuffer Buffer, ulong Size) binding3) + { + var buffers = stackalloc DescriptorBufferInfo[4] + { + new DescriptorBufferInfo { Buffer = binding0.Buffer, Offset = 0, Range = binding0.Size }, + new DescriptorBufferInfo { Buffer = binding1.Buffer, Offset = 0, Range = binding1.Size }, + new DescriptorBufferInfo { Buffer = binding2.Buffer, Offset = 0, Range = binding2.Size }, + new DescriptorBufferInfo { Buffer = binding3.Buffer, Offset = 0, Range = binding3.Size }, + }; + + var writes = stackalloc WriteDescriptorSet[4]; + for (uint index = 0; index < 4; index++) + { + writes[index] = new WriteDescriptorSet + { + SType = StructureType.WriteDescriptorSet, + DstSet = descriptorSet, + DstBinding = index, + DstArrayElement = 0, + DescriptorCount = 1, + DescriptorType = DescriptorType.StorageBuffer, + PBufferInfo = &buffers[index], + }; + } + + _vk.UpdateDescriptorSets(_device, 4, writes, 0, null); + } + + private CommandBuffer AllocateCommandBuffer() + { + var allocateInfo = new CommandBufferAllocateInfo + { + SType = StructureType.CommandBufferAllocateInfo, + CommandPool = _commandPool, + Level = CommandBufferLevel.Primary, + CommandBufferCount = 1, + }; + Check( + _vk.AllocateCommandBuffers(_device, &allocateInfo, out var commandBuffer), + "vkAllocateCommandBuffers(detile)"); + return commandBuffer; + } + + private void BeginCommandBuffer(CommandBuffer commandBuffer) + { + var beginInfo = new CommandBufferBeginInfo + { + SType = StructureType.CommandBufferBeginInfo, + Flags = CommandBufferUsageFlags.OneTimeSubmitBit, + }; + Check(_vk.BeginCommandBuffer(commandBuffer, &beginInfo), "vkBeginCommandBuffer(detile)"); + } + + private void TransitionImage( + CommandBuffer commandBuffer, + Image image, + ImageLayout oldLayout, + ImageLayout newLayout, + AccessFlags srcAccess, + AccessFlags dstAccess, + PipelineStageFlags srcStage, + PipelineStageFlags dstStage, + uint layers) + { + var barrier = new ImageMemoryBarrier + { + SType = StructureType.ImageMemoryBarrier, + SrcAccessMask = srcAccess, + DstAccessMask = dstAccess, + OldLayout = oldLayout, + NewLayout = newLayout, + SrcQueueFamilyIndex = Vk.QueueFamilyIgnored, + DstQueueFamilyIndex = Vk.QueueFamilyIgnored, + Image = image, + SubresourceRange = new ImageSubresourceRange(ImageAspectFlags.ColorBit, 0, 1, 0, layers), + }; + _vk.CmdPipelineBarrier(commandBuffer, srcStage, dstStage, 0, 0, null, 0, null, 1, &barrier); + } + + private Fence CreateFence() + { + var fenceInfo = new FenceCreateInfo { SType = StructureType.FenceCreateInfo }; + Check(_vk.CreateFence(_device, &fenceInfo, null, out var fence), "vkCreateFence(detile)"); + return fence; + } + + private void DestroyBuffer(VkBuffer buffer, DeviceMemory memory) + { + if (buffer.Handle != 0) + { + _vk.DestroyBuffer(_device, buffer, null); + } + + if (memory.Handle != 0) + { + _vk.FreeMemory(_device, memory, null); + } + } + + private void Check(Result result, string operation) + { + if (result != Result.Success) + { + throw new InvalidOperationException($"{operation} failed: {result}"); + } + } + + public void Dispose() + { + if (_disposed) + { + return; + } + + _disposed = true; + if (_pipeline.Handle != 0) + { + _vk.DestroyPipeline(_device, _pipeline, null); + } + + if (_pipelineLayout.Handle != 0) + { + _vk.DestroyPipelineLayout(_device, _pipelineLayout, null); + } + + if (_descriptorSetLayout.Handle != 0) + { + _vk.DestroyDescriptorSetLayout(_device, _descriptorSetLayout, null); + } + + if (_shaderModule.Handle != 0) + { + _vk.DestroyShaderModule(_device, _shaderModule, null); + } + + if (_commandPool.Handle != 0) + { + _vk.DestroyCommandPool(_device, _commandPool, null); + } + } +} diff --git a/src/SharpEmu.Libs/VideoOut/VulkanDetileSelfTest.cs b/src/SharpEmu.Libs/VideoOut/VulkanDetileSelfTest.cs new file mode 100644 index 00000000..cd699577 --- /dev/null +++ b/src/SharpEmu.Libs/VideoOut/VulkanDetileSelfTest.cs @@ -0,0 +1,498 @@ +// Copyright (C) 2026 SharpEmu Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +using SharpEmu.Libs.Agc; +using Silk.NET.Vulkan; +using VkBuffer = Silk.NET.Vulkan.Buffer; + +namespace SharpEmu.Libs.VideoOut; + +/// +/// Opt-in GPU equivalence check for . When +/// SHARPEMU_DETILE_SELFTEST=1 it builds a known tiled surface, deswizzles it on +/// the GPU into a real image, reads the image back, and compares against the CPU +/// — the same equivalence the unit test proves +/// for the params, now end-to-end through the actual Vulkan pass. It logs +/// [DETILE-SELFTEST] PASS/FAIL and never throws into startup (any failure is +/// caught and logged), so it is safe to leave wired. +/// +internal static unsafe class VulkanDetileSelfTest +{ + private const uint Width = 256; + private const uint Height = 256; + + // (swizzleMode, bytesPerElement, image format). Mode 27 is exact-XOR, mode 8 + // (64 KiB Z) is block-table — both branches. bpp 4/8/16 exercises the + // one/two/four-words-per-element copy. These formats are non-block-compressed + // (element grid == texel grid), so Width/Height are both element and texel dims. + private static readonly (uint Mode, int Bpp, Format Format)[] Cases = + [ + (27, 4, Format.R8G8B8A8Unorm), + (8, 4, Format.R8G8B8A8Unorm), + (27, 8, Format.R32G32Uint), + (27, 16, Format.R32G32B32A32Uint), + ]; + + public static void RunIfRequested( + Vk vk, + Device device, + Queue queue, + PhysicalDevice physicalDevice, + uint queueFamilyIndex) + { + if (Environment.GetEnvironmentVariable("SHARPEMU_DETILE_SELFTEST") != "1") + { + return; + } + + try + { + Run(vk, device, queue, physicalDevice, queueFamilyIndex); + } + catch (Exception exception) + { + Console.Error.WriteLine($"[DETILE-SELFTEST] FAIL (exception): {exception.Message}"); + } + } + + private static void Run( + Vk vk, + Device device, + Queue queue, + PhysicalDevice physicalDevice, + uint queueFamilyIndex) + { + using var pass = new VulkanDetilePass(vk, device, queue, physicalDevice, queueFamilyIndex); + var commandPool = CreateCommandPool(vk, device, queueFamilyIndex); + try + { + foreach (var (mode, bpp, format) in Cases) + { + // A plain 2D texture (1 layer) and an array texture (2 layers) — the + // arrayed case exercises the kernel's dispatch-Z slice addressing. + RunCase(vk, device, physicalDevice, queue, commandPool, pass, mode, bpp, format, layers: 1); + RunCase(vk, device, physicalDevice, queue, commandPool, pass, mode, bpp, format, layers: 2); + } + } + finally + { + if (commandPool.Handle != 0) + { + vk.DestroyCommandPool(device, commandPool, null); + } + } + } + + private static void RunCase( + Vk vk, + Device device, + PhysicalDevice physicalDevice, + Queue queue, + CommandPool commandPool, + VulkanDetilePass pass, + uint swizzleMode, + int bytesPerElement, + Format format, + uint layers) + { + var parameters = GnmTiling.GetDetileParams(swizzleMode, bytesPerElement, (int)Width, (int)Height); + if (!parameters.IsSupported || !VulkanDetilePass.Supports(parameters)) + { + Console.Error.WriteLine( + $"[DETILE-SELFTEST] FAIL: mode {swizzleMode} bpp {bytesPerElement} not supported by the GPU pass."); + return; + } + + // Whole-block tiled source with a per-layer-distinct deterministic pattern + // (so a slice mix-up is caught), the array slices packed contiguously. + var blocksHigh = ((int)Height + parameters.BlockHeight - 1) / parameters.BlockHeight; + var sliceTiledBytes = (int)((long)parameters.BlocksPerRow * blocksHigh * parameters.BlockBytes); + var sliceLinearBytes = (int)(Width * Height * bytesPerElement); + var tiled = new byte[sliceTiledBytes * layers]; + var expected = new byte[sliceLinearBytes * layers]; + for (var layer = 0; layer < layers; layer++) + { + for (var index = 0; index < sliceTiledBytes; index++) + { + tiled[layer * sliceTiledBytes + index] = (byte)((index * 31 + 7 + layer * 101) & 0xFF); + } + + if (!GnmTiling.TryDetile( + tiled.AsSpan(layer * sliceTiledBytes, sliceTiledBytes), + expected.AsSpan(layer * sliceLinearBytes, sliceLinearBytes), + swizzleMode, (int)Width, (int)Height, bytesPerElement)) + { + Console.Error.WriteLine("[DETILE-SELFTEST] FAIL: CPU TryDetile declined."); + return; + } + } + + var tiledBytes = tiled; + var label = $"mode{swizzleMode} {bytesPerElement}bpp x{layers}"; + + // Phase 1: the one-shot DetileIntoImage (submit + wait in place). + VerifyPhase( + vk, device, physicalDevice, queue, commandPool, expected, layers, format, $"DetileIntoImage {label}", + image => pass.DetileIntoImage(image, ImageLayout.Undefined, Width, Height, layers, tiledBytes, parameters)); + + // Phase 2: RecordDetile — the exact code path the render loop uses + // (record into a command buffer, submit, retire the transients). + VerifyPhase( + vk, device, physicalDevice, queue, commandPool, expected, layers, format, $"RecordDetile {label}", + image => RecordDetileAndSubmit(vk, device, queue, commandPool, pass, image, layers, tiledBytes, parameters)); + } + + private static void VerifyPhase( + Vk vk, + Device device, + PhysicalDevice physicalDevice, + Queue queue, + CommandPool commandPool, + byte[] expected, + uint layers, + Format format, + string label, + Func detile) + { + var image = CreateImage(vk, device, physicalDevice, format, layers, out var imageMemory); + var readback = CreateHostBuffer( + vk, device, physicalDevice, (ulong)expected.Length, BufferUsageFlags.TransferDstBit, out var readbackMemory); + try + { + if (!detile(image)) + { + Console.Error.WriteLine($"[DETILE-SELFTEST] {label} FAIL: declined."); + return; + } + + CopyImageToBuffer(vk, device, queue, commandPool, image, readback, layers); + + void* mapped; + Check( + vk.MapMemory(device, readbackMemory, 0, (ulong)expected.Length, 0, &mapped), + $"vkMapMemory(selftest {label})"); + var actual = new Span(mapped, expected.Length); + var firstMismatch = -1; + for (var index = 0; index < expected.Length; index++) + { + if (actual[index] != expected[index]) + { + firstMismatch = index; + break; + } + } + + vk.UnmapMemory(device, readbackMemory); + + Console.Error.WriteLine(firstMismatch < 0 + ? $"[DETILE-SELFTEST] {label} PASS: {Width}x{Height}x{layers} matches CPU detile ({expected.Length} bytes)." + : $"[DETILE-SELFTEST] {label} FAIL: first mismatch at byte {firstMismatch}."); + } + finally + { + if (readback.Handle != 0) + { + vk.DestroyBuffer(device, readback, null); + } + + if (readbackMemory.Handle != 0) + { + vk.FreeMemory(device, readbackMemory, null); + } + + if (image.Handle != 0) + { + vk.DestroyImage(device, image, null); + } + + if (imageMemory.Handle != 0) + { + vk.FreeMemory(device, imageMemory, null); + } + } + } + + // Records the detile into a fresh command buffer, submits, waits, and retires + // the transients exactly as the presenter's batch does — verifying the render + // path's code (RecordDetile) without needing a game to trigger it. + private static bool RecordDetileAndSubmit( + Vk vk, + Device device, + Queue queue, + CommandPool commandPool, + VulkanDetilePass pass, + Image image, + uint layers, + ReadOnlySpan tiled, + in DetileParams parameters) + { + var commandBuffer = AllocateCommandBuffer(vk, device, commandPool); + var beginInfo = new CommandBufferBeginInfo + { + SType = StructureType.CommandBufferBeginInfo, + Flags = CommandBufferUsageFlags.OneTimeSubmitBit, + }; + Check(vk.BeginCommandBuffer(commandBuffer, &beginInfo), "vkBeginCommandBuffer(selftest record)"); + + if (!pass.RecordDetile( + commandBuffer, image, ImageLayout.Undefined, Width, Height, layers, tiled, parameters, out var transients)) + { + _ = vk.EndCommandBuffer(commandBuffer); + vk.FreeCommandBuffers(device, commandPool, 1, &commandBuffer); + return false; + } + + Check(vk.EndCommandBuffer(commandBuffer), "vkEndCommandBuffer(selftest record)"); + + var fenceInfo = new FenceCreateInfo { SType = StructureType.FenceCreateInfo }; + Check(vk.CreateFence(device, &fenceInfo, null, out var fence), "vkCreateFence(selftest record)"); + try + { + var submitInfo = new SubmitInfo + { + SType = StructureType.SubmitInfo, + CommandBufferCount = 1, + PCommandBuffers = &commandBuffer, + }; + Check(vk.QueueSubmit(queue, 1, &submitInfo, fence), "vkQueueSubmit(selftest record)"); + Check(vk.WaitForFences(device, 1, &fence, true, ulong.MaxValue), "vkWaitForFences(selftest record)"); + } + finally + { + vk.DestroyFence(device, fence, null); + vk.FreeCommandBuffers(device, commandPool, 1, &commandBuffer); + foreach (var (buffer, memory) in transients.Buffers) + { + if (buffer.Handle != 0) + { + vk.DestroyBuffer(device, buffer, null); + } + + if (memory.Handle != 0) + { + vk.FreeMemory(device, memory, null); + } + } + + if (transients.DescriptorPool.Handle != 0) + { + vk.DestroyDescriptorPool(device, transients.DescriptorPool, null); + } + } + + return true; + } + + private static Image CreateImage( + Vk vk, + Device device, + PhysicalDevice physicalDevice, + Format format, + uint layers, + out DeviceMemory memory) + { + var imageInfo = new ImageCreateInfo + { + SType = StructureType.ImageCreateInfo, + ImageType = ImageType.Type2D, + Format = format, + Extent = new Extent3D(Width, Height, 1), + MipLevels = 1, + ArrayLayers = layers, + Samples = SampleCountFlags.Count1Bit, + Tiling = ImageTiling.Optimal, + Usage = ImageUsageFlags.TransferDstBit | ImageUsageFlags.TransferSrcBit, + SharingMode = SharingMode.Exclusive, + InitialLayout = ImageLayout.Undefined, + }; + Check(vk.CreateImage(device, &imageInfo, null, out var image), "vkCreateImage(selftest)"); + + vk.GetImageMemoryRequirements(device, image, out var requirements); + var allocateInfo = new MemoryAllocateInfo + { + SType = StructureType.MemoryAllocateInfo, + AllocationSize = requirements.Size, + MemoryTypeIndex = FindMemoryType( + vk, + physicalDevice, + requirements.MemoryTypeBits, + MemoryPropertyFlags.DeviceLocalBit), + }; + Check(vk.AllocateMemory(device, &allocateInfo, null, out memory), "vkAllocateMemory(selftest image)"); + Check(vk.BindImageMemory(device, image, memory, 0), "vkBindImageMemory(selftest)"); + return image; + } + + private static void CopyImageToBuffer( + Vk vk, + Device device, + Queue queue, + CommandPool commandPool, + Image image, + VkBuffer destination, + uint layers) + { + var commandBuffer = AllocateCommandBuffer(vk, device, commandPool); + var beginInfo = new CommandBufferBeginInfo + { + SType = StructureType.CommandBufferBeginInfo, + Flags = CommandBufferUsageFlags.OneTimeSubmitBit, + }; + Check(vk.BeginCommandBuffer(commandBuffer, &beginInfo), "vkBeginCommandBuffer(selftest readback)"); + + // DetileIntoImage left the image ShaderReadOnly; move it to TransferSrc. + var toTransferSrc = new ImageMemoryBarrier + { + SType = StructureType.ImageMemoryBarrier, + SrcAccessMask = AccessFlags.ShaderReadBit, + DstAccessMask = AccessFlags.TransferReadBit, + OldLayout = ImageLayout.ShaderReadOnlyOptimal, + NewLayout = ImageLayout.TransferSrcOptimal, + SrcQueueFamilyIndex = Vk.QueueFamilyIgnored, + DstQueueFamilyIndex = Vk.QueueFamilyIgnored, + Image = image, + SubresourceRange = new ImageSubresourceRange(ImageAspectFlags.ColorBit, 0, 1, 0, layers), + }; + vk.CmdPipelineBarrier( + commandBuffer, + PipelineStageFlags.FragmentShaderBit, + PipelineStageFlags.TransferBit, + 0, + 0, + null, + 0, + null, + 1, + &toTransferSrc); + + // Layer-major readback: one copy pulls every array slice back into the + // buffer contiguously, matching the packed `expected` layout. + var region = new BufferImageCopy + { + BufferOffset = 0, + BufferRowLength = 0, + BufferImageHeight = 0, + ImageSubresource = new ImageSubresourceLayers(ImageAspectFlags.ColorBit, 0, 0, layers), + ImageOffset = default, + ImageExtent = new Extent3D(Width, Height, 1), + }; + vk.CmdCopyImageToBuffer( + commandBuffer, + image, + ImageLayout.TransferSrcOptimal, + destination, + 1, + ®ion); + + Check(vk.EndCommandBuffer(commandBuffer), "vkEndCommandBuffer(selftest readback)"); + + var fenceInfo = new FenceCreateInfo { SType = StructureType.FenceCreateInfo }; + Check(vk.CreateFence(device, &fenceInfo, null, out var fence), "vkCreateFence(selftest)"); + try + { + var submitInfo = new SubmitInfo + { + SType = StructureType.SubmitInfo, + CommandBufferCount = 1, + PCommandBuffers = &commandBuffer, + }; + Check(vk.QueueSubmit(queue, 1, &submitInfo, fence), "vkQueueSubmit(selftest readback)"); + Check( + vk.WaitForFences(device, 1, &fence, true, ulong.MaxValue), + "vkWaitForFences(selftest readback)"); + } + finally + { + vk.DestroyFence(device, fence, null); + vk.FreeCommandBuffers(device, commandPool, 1, &commandBuffer); + } + } + + private static CommandPool CreateCommandPool(Vk vk, Device device, uint queueFamilyIndex) + { + var poolInfo = new CommandPoolCreateInfo + { + SType = StructureType.CommandPoolCreateInfo, + QueueFamilyIndex = queueFamilyIndex, + Flags = CommandPoolCreateFlags.ResetCommandBufferBit, + }; + Check(vk.CreateCommandPool(device, &poolInfo, null, out var pool), "vkCreateCommandPool(selftest)"); + return pool; + } + + private static CommandBuffer AllocateCommandBuffer(Vk vk, Device device, CommandPool commandPool) + { + var allocateInfo = new CommandBufferAllocateInfo + { + SType = StructureType.CommandBufferAllocateInfo, + CommandPool = commandPool, + Level = CommandBufferLevel.Primary, + CommandBufferCount = 1, + }; + Check( + vk.AllocateCommandBuffers(device, &allocateInfo, out var commandBuffer), + "vkAllocateCommandBuffers(selftest)"); + return commandBuffer; + } + + private static VkBuffer CreateHostBuffer( + Vk vk, + Device device, + PhysicalDevice physicalDevice, + ulong size, + BufferUsageFlags usage, + out DeviceMemory memory) + { + var bufferInfo = new BufferCreateInfo + { + SType = StructureType.BufferCreateInfo, + Size = size, + Usage = usage, + SharingMode = SharingMode.Exclusive, + }; + Check(vk.CreateBuffer(device, &bufferInfo, null, out var buffer), "vkCreateBuffer(selftest)"); + + vk.GetBufferMemoryRequirements(device, buffer, out var requirements); + var allocateInfo = new MemoryAllocateInfo + { + SType = StructureType.MemoryAllocateInfo, + AllocationSize = requirements.Size, + MemoryTypeIndex = FindMemoryType( + vk, + physicalDevice, + requirements.MemoryTypeBits, + MemoryPropertyFlags.HostVisibleBit | MemoryPropertyFlags.HostCoherentBit), + }; + Check(vk.AllocateMemory(device, &allocateInfo, null, out memory), "vkAllocateMemory(selftest buffer)"); + Check(vk.BindBufferMemory(device, buffer, memory, 0), "vkBindBufferMemory(selftest)"); + return buffer; + } + + private static uint FindMemoryType( + Vk vk, + PhysicalDevice physicalDevice, + uint typeBits, + MemoryPropertyFlags requiredFlags) + { + vk.GetPhysicalDeviceMemoryProperties(physicalDevice, out var properties); + var memoryTypes = &properties.MemoryTypes.Element0; + for (uint index = 0; index < properties.MemoryTypeCount; index++) + { + if ((typeBits & (1u << (int)index)) != 0 && + (memoryTypes[index].PropertyFlags & requiredFlags) == requiredFlags) + { + return index; + } + } + + throw new InvalidOperationException("No compatible Vulkan memory type for the detile self-test."); + } + + private static void Check(Result result, string operation) + { + if (result != Result.Success) + { + throw new InvalidOperationException($"{operation} failed: {result}"); + } + } +} diff --git a/src/SharpEmu.Libs/VideoOut/VulkanVideoPresenter.cs b/src/SharpEmu.Libs/VideoOut/VulkanVideoPresenter.cs index e6a9088a..1ecbb450 100644 --- a/src/SharpEmu.Libs/VideoOut/VulkanVideoPresenter.cs +++ b/src/SharpEmu.Libs/VideoOut/VulkanVideoPresenter.cs @@ -2860,6 +2860,14 @@ internal static unsafe class VulkanVideoPresenter private long _lastPipelineCacheSaveTick; private Queue _queue; private uint _queueFamilyIndex; + // GPU deswizzle (default on; SHARPEMU_GPU_DETILE=0 forces the CPU path). + // Lazily built on the first tiled texture; disposed with the presenter. + private static readonly bool _gpuDetileEnabled = !string.Equals( + Environment.GetEnvironmentVariable("SHARPEMU_GPU_DETILE"), "0", StringComparison.Ordinal); + private static readonly bool _gpuDetileLog = string.Equals( + Environment.GetEnvironmentVariable("SHARPEMU_LOG_GPU_DETILE"), "1", StringComparison.Ordinal); + private VulkanDetilePass? _detilePass; + private long _gpuDetileCount; private SwapchainKHR _swapchain; private Image[] _swapchainImages = []; private ImageView[] _swapchainImageViews = []; @@ -2906,6 +2914,9 @@ internal static unsafe class VulkanVideoPresenter private readonly Stack _recycledGuestFences = new(); private readonly Stack _recycledGuestCommandBuffers = new(); private readonly List<(VkBuffer Buffer, DeviceMemory Memory)> _batchRetireBuffers = new(); + // Descriptor pools from GPU-detile passes recorded into the batch; retired + // with the batch's fence, alongside _batchRetireBuffers. + private readonly List _batchRetireDescriptorPools = new(); private const int MaxRecycledGuestFences = 32; private const int MaxRecycledGuestCommandBuffers = 32; private VkBuffer _stagingBuffer; @@ -3277,6 +3288,7 @@ internal static unsafe class VulkanVideoPresenter IReadOnlyList Resources, IReadOnlyList TraceImages, IReadOnlyList<(VkBuffer Buffer, DeviceMemory Memory)> RetireBuffers, + IReadOnlyList RetirePools, ulong Timeline, string DebugName, VulkanGuestQueueIdentity Queue, @@ -4241,6 +4253,7 @@ internal static unsafe class VulkanVideoPresenter _vk.GetDeviceQueue(_device, _queueFamilyIndex, 0, out _queue); LoadDebugUtilsCommands(); + VulkanDetileSelfTest.RunIfRequested(_vk, _device, _queue, _physicalDevice, _queueFamilyIndex); if (!_vk.TryGetDeviceExtension(_instance, _device, out _swapchainApi)) { throw new InvalidOperationException("VK_KHR_swapchain is unavailable."); @@ -4949,7 +4962,10 @@ internal static unsafe class VulkanVideoPresenter _batchCommandBuffer, _batchResources.ToArray(), _batchTraceImages.ToArray(), - _batchRetireBuffers.Count > 0 ? _batchRetireBuffers.ToArray() : []); + _batchRetireBuffers.Count > 0 ? _batchRetireBuffers.ToArray() : [], + retirePools: _batchRetireDescriptorPools.Count > 0 + ? _batchRetireDescriptorPools.ToArray() + : []); } catch { @@ -4967,6 +4983,11 @@ internal static unsafe class VulkanVideoPresenter _vk.FreeMemory(_device, memory, null); } + foreach (var pool in _batchRetireDescriptorPools) + { + _vk.DestroyDescriptorPool(_device, pool, null); + } + ReleaseGuestCommandBuffer(_batchCommandBuffer); throw; } @@ -4975,6 +4996,7 @@ internal static unsafe class VulkanVideoPresenter _batchResources.Clear(); _batchTraceImages.Clear(); _batchRetireBuffers.Clear(); + _batchRetireDescriptorPools.Clear(); _batchCommandBuffer = default; } } @@ -4984,7 +5006,8 @@ internal static unsafe class VulkanVideoPresenter IReadOnlyList resources, IReadOnlyList traceImages, IReadOnlyList<(VkBuffer Buffer, DeviceMemory Memory)>? retireBuffers = null, - IReadOnlyList? referencedResources = null) + IReadOnlyList? referencedResources = null, + IReadOnlyList? retirePools = null) { var fence = AcquireGuestFence(); try @@ -5042,6 +5065,7 @@ internal static unsafe class VulkanVideoPresenter resources, traceImages, retireBuffers ?? [], + retirePools ?? [], _submitTimeline, resources.Count > 0 ? resources[0].DebugName : "batch", _activeGuestQueue, @@ -5213,6 +5237,11 @@ internal static unsafe class VulkanVideoPresenter _vk.FreeMemory(_device, memory, null); } + foreach (var pool in submission.RetirePools) + { + _vk.DestroyDescriptorPool(_device, pool, null); + } + ReleaseGuestCommandBuffer(submission.CommandBuffer); ReleaseGuestFence(submission.Fence, needsReset: true); if (submission.Timeline > _completedTimeline) @@ -7845,7 +7874,7 @@ internal static unsafe class VulkanVideoPresenter MarkTextureContentCached(key); SharpEmu.HLE.GuestImageWriteTracker.Track( texture.Address, - (ulong)texture.RgbaPixels.Length, + (ulong)(texture.TiledSource?.Length ?? texture.RgbaPixels.Length), CurrentGuestWorkSequenceForDiagnostics, "vulkan.texture-cache"); } @@ -8226,6 +8255,10 @@ internal static unsafe class VulkanVideoPresenter return (uint)selectedMipLevel; } + private VulkanDetilePass EnsureDetilePass() => + _detilePass ??= new VulkanDetilePass( + _vk, _device, _queue, _physicalDevice, _queueFamilyIndex); + private TextureResource CreateTextureResource(GuestDrawTexture texture) { var width = Math.Max(texture.Width, 1); @@ -8255,31 +8288,91 @@ internal static unsafe class VulkanVideoPresenter $"dst=0x{texture.DstSelect:X3} " + $"bytes={texture.RgbaPixels.Length} expected={expectedSize}"); } - var pixels = texture.RgbaPixels.Length == (int)(expectedSize * layers) - ? texture.RgbaPixels - : CreateFallbackTexturePixels(texture.Format, rowLength, height, expectedSize); - if (!ReferenceEquals(pixels, texture.RgbaPixels)) + // The GPU detile pass deswizzles plain 2D and array textures (one + // dispatch-Z layer per slice) at 4/8/16 bpp, including block-compressed + // formats (element grid = ceil(texels/4), smaller than the texel grid). + // Validate against the element grid + bpp from the resolved params, and + // require the tiled source to cover every layer's linear extent (tiled + // slices are >= the linear size due to whole-block padding). + DetileParams? gpuDetileParams = null; + byte[]? gpuTiledSource = null; + if (_gpuDetileEnabled && + texture.Detile is { } detileCandidate && + texture.TiledSource is { Length: > 0 } tiledCandidate && + VulkanDetilePass.Supports(detileCandidate) && + detileCandidate.ElementsWide > 0 && + detileCandidate.ElementsHigh > 0 && + (long)tiledCandidate.Length >= + (long)detileCandidate.ElementsWide * detileCandidate.ElementsHigh * + detileCandidate.BytesPerElement * layers && + tiledCandidate.Length % (int)(layers * (uint)detileCandidate.BytesPerElement) == 0) { - layers = 1; + gpuDetileParams = detileCandidate; + gpuTiledSource = tiledCandidate; } - if (AddressListContains("SHARPEMU_FORCE_WHITE_TEXTURE_TARGETS", texture.Address)) - { - pixels = pixels.ToArray(); - pixels.AsSpan().Fill(0xFF); - Console.Error.WriteLine( - $"[LOADER][TRACE] vk.texture_force_white addr=0x{texture.Address:X16} " + - $"size={width}x{height} bytes={pixels.Length}"); - } - DumpTextureUpload(texture, pixels, rowLength, width, height); - TraceTextureUploadContents(texture, pixels, rowLength, width, height, vkFormat); - var uploadPixels = texture.Format == 13 - ? ExpandRgb32Pixels(pixels) - : pixels; - var contentFingerprint = ComputeTextureContentFingerprint(pixels); - var (stagingBuffer, stagingMemory) = CreateTextureStagingBuffer( - uploadPixels, - $"{TextureDebugName(texture, vkFormat)} staging"); + VkBuffer stagingBuffer = default; + DeviceMemory stagingMemory = default; + ulong contentFingerprint; + if (gpuTiledSource is { } gpuSource) + { + // GPU detile: no CPU staging; the compute pass writes the image directly. + contentFingerprint = ComputeTextureContentFingerprint(gpuSource); + } + else + { + // Safety net: the AGC gate can package a texture as a GPU-detile + // candidate (empty RgbaPixels + TiledSource) that this path did + // not accept for the GPU compute pass (see the gpuTiledSource + // guard above). Detile the raw tiled bytes on the CPU here rather + // than letting empty RgbaPixels fall through to a blank fallback + // image — otherwise such textures render empty (missing text). + var cpuDetiled = texture.RgbaPixels; + if (cpuDetiled.Length == 0 && + layers == 1 && + texture.TiledSource is { Length: > 0 } fallbackTiled && + texture.Detile is { } fallbackParams && + expectedSize > 0 && + expectedSize <= int.MaxValue) + { + var linear = new byte[expectedSize]; + if (GnmTiling.TryDetile( + fallbackTiled, + linear, + texture.TileMode, + fallbackParams.ElementsWide, + fallbackParams.ElementsHigh, + fallbackParams.BytesPerElement)) + { + cpuDetiled = linear; + } + } + + var pixels = cpuDetiled.Length == (int)(expectedSize * layers) + ? cpuDetiled + : CreateFallbackTexturePixels(texture.Format, rowLength, height, expectedSize); + if (!ReferenceEquals(pixels, texture.RgbaPixels)) + { + layers = 1; + } + if (AddressListContains("SHARPEMU_FORCE_WHITE_TEXTURE_TARGETS", texture.Address)) + { + pixels = pixels.ToArray(); + pixels.AsSpan().Fill(0xFF); + Console.Error.WriteLine( + $"[LOADER][TRACE] vk.texture_force_white addr=0x{texture.Address:X16} " + + $"size={width}x{height} bytes={pixels.Length}"); + } + DumpTextureUpload(texture, pixels, rowLength, width, height); + TraceTextureUploadContents(texture, pixels, rowLength, width, height, vkFormat); + var uploadPixels = texture.Format == 13 + ? ExpandRgb32Pixels(pixels) + : pixels; + contentFingerprint = ComputeTextureContentFingerprint(pixels); + (stagingBuffer, stagingMemory) = CreateTextureStagingBuffer( + uploadPixels, + $"{TextureDebugName(texture, vkFormat)} staging"); + } var supportsMutableUsage = !IsBlockCompressedFormat(vkFormat); var supportsAttachmentUsage = @@ -8341,6 +8434,83 @@ internal static unsafe class VulkanVideoPresenter var debugName = TextureDebugName(texture, vkFormat); SetDebugName(ObjectType.Image, image.Handle, $"{debugName} image"); SetDebugName(ObjectType.ImageView, view.Handle, $"{debugName} view"); + + // GPU detile: record the deswizzle into the shared batch command buffer + // (async — never a blocking submit on the render thread), leaving the + // image ShaderReadOnly before the draw that samples it. The transient + // buffers + descriptor pool retire with the batch fence. On any failure + // fall back to a CPU detile + normal staged upload. + var gpuDetiled = false; + if (gpuTiledSource is { } detileSource && gpuDetileParams is { } detileParameters) + { + try + { + var detileCommandBuffer = BeginBatchedGuestCommands(); + CloseOpenTranslatedRenderPass(); + if (EnsureDetilePass().RecordDetile( + detileCommandBuffer, + image, + ImageLayout.Undefined, + width, + height, + layers, + detileSource, + detileParameters, + out var detileTransients)) + { + _batchRetireBuffers.AddRange(detileTransients.Buffers); + _batchRetireDescriptorPools.Add(detileTransients.DescriptorPool); + gpuDetiled = true; + } + } + catch (Exception exception) + { + Console.Error.WriteLine( + $"[LOADER][WARN] GPU detile failed for addr=0x{texture.Address:X16}, " + + $"falling back to CPU: {exception.Message}"); + gpuDetiled = false; + } + + if (!gpuDetiled) + { + // CPU fallback: the tiled source packs the array slices + // contiguously, so detile each slice into its layer-major + // linear region (single layer degrades to one iteration). + var totalLinear = checked((int)(expectedSize * layers)); + var linear = new byte[totalLinear]; + var sliceTiledBytes = detileSource.Length / (int)layers; + var sliceLinearBytes = (int)expectedSize; + var detiledAll = true; + for (var layer = 0; layer < layers; layer++) + { + // TryDetile iterates the element grid (for BC, ceil(texels/4)). + if (!GnmTiling.TryDetile( + detileSource.AsSpan(layer * sliceTiledBytes, sliceTiledBytes), + linear.AsSpan(layer * sliceLinearBytes, sliceLinearBytes), + texture.TileMode, + detileParameters.ElementsWide, + detileParameters.ElementsHigh, + detileParameters.BytesPerElement)) + { + detiledAll = false; + break; + } + } + + if (detiledAll) + { + (stagingBuffer, stagingMemory) = CreateTextureStagingBuffer( + linear, $"{TextureDebugName(texture, vkFormat)} staging(cpu-fallback)"); + } + } + else if (_gpuDetileLog && Interlocked.Increment(ref _gpuDetileCount) is 1 or 100 or 1000 or 10000) + { + Console.Error.WriteLine( + $"[GPU-DETILE] active: {_gpuDetileCount} texture(s) detiled on GPU " + + $"(latest {width}x{height} mode {texture.TileMode})."); + } + } + var resource = new TextureResource { Address = texture.Address, @@ -8356,7 +8526,7 @@ internal static unsafe class VulkanVideoPresenter RowLength = rowLength, DstSelect = texture.DstSelect, Layers = layers, - NeedsUpload = true, + NeedsUpload = !gpuDetiled, OwnsStorage = true, SamplerState = texture.Sampler, CpuContentFingerprint = contentFingerprint, @@ -8367,6 +8537,7 @@ internal static unsafe class VulkanVideoPresenter if (texture.Address != 0 && !texture.ArrayedView && layers == 1 && + !gpuDetiled && !_guestImages.ContainsKey(texture.Address)) { var guestFormat = GetGuestTextureFormat(texture.Format, texture.NumberType); @@ -16643,6 +16814,8 @@ internal static unsafe class VulkanVideoPresenter _lastOrderedGuestFlipVersions.Clear(); } DestroySwapchainResources(); + _detilePass?.Dispose(); + _detilePass = null; if (_device.Handle != 0) { if (_pipelineCache.Handle != 0) diff --git a/src/SharpEmu.ShaderCompiler.Metal/MslFixedShaders.cs b/src/SharpEmu.ShaderCompiler.Metal/MslFixedShaders.cs index ab301943..2cccb7de 100644 --- a/src/SharpEmu.ShaderCompiler.Metal/MslFixedShaders.cs +++ b/src/SharpEmu.ShaderCompiler.Metal/MslFixedShaders.cs @@ -76,6 +76,14 @@ public static class MslFixedShaders /// public static string CreateDepthOnlyFragment() => MslTemplates.Render("depth_only_fragment"); + /// + /// Compute kernel that deswizzles one RDNA2 exact-XOR tiled surface (swizzle + /// modes 5/9/24/27) at 4 bytes/element into a linear output buffer — the MSL + /// twin of SpirvFixedShaders.CreateDetileCompute. Entry point + /// "detile_cs"; buffers 0=tiled, 1=xTerm, 2=yTerm, 3=out, 4=DetileParams. + /// + public static string CreateDetileCompute() => MslTemplates.Render("detile_compute"); + private static string Format(float value) => value.ToString("0.0######", CultureInfo.InvariantCulture) + "f"; } diff --git a/src/SharpEmu.ShaderCompiler.Metal/Templates/detile_compute.msl b/src/SharpEmu.ShaderCompiler.Metal/Templates/detile_compute.msl new file mode 100644 index 00000000..6b077fd7 --- /dev/null +++ b/src/SharpEmu.ShaderCompiler.Metal/Templates/detile_compute.msl @@ -0,0 +1,70 @@ +#include + +using namespace metal; + +// GPU deswizzle for RDNA2 tiled surfaces at 4/8/16 bytes/element — the MSL twin of +// SpirvFixedShaders.CreateDetileCompute and a direct mirror of +// GnmTiling.GetDetileParams. Handles both supported equation families and both +// plain 2D textures and array textures (one grid-Z layer per slice; the caller +// packs the tiled slices contiguously). width/height are ELEMENT dims (for +// block-compressed formats a 4x4 block is one element); each element spans +// uintsPerElement = bpp/4 words, and the X grid is widened by that factor so each +// thread copies one word (elemX = gidX / upe, word = gidX % upe): +// inBlock = equation == 1 // BlockTable, modes 1/4/8 +// ? blockTable[(y % blockHeight) * blockWidth + (elemX % blockWidth)] +// : xTerm[elemX & xMask] ^ yTerm[y & yMask]; // ExactXor 5/9/24/27 +// srcElem = z * srcSliceElements +// + (y / blockHeight * blocksPerRow + elemX / blockWidth) * blockElements +// + inBlock; +// dstElem = z * width * height + y * width + elemX; +// out[dstElem * upe + word] = tiled[srcElem * upe + word]; +// Term tables hold ELEMENT offsets. buffer(1) carries xTerm (ExactXor) OR the +// block table (BlockTable); the two equations index different-sized buffers, so +// exactly one branch runs. 1/2 bpp are sub-word and stay on the CPU. +struct DetileParams +{ + uint width; + uint height; + uint blockWidth; + uint blockHeight; + uint blockElements; + uint blocksPerRow; + uint xMask; + uint yMask; + uint srcSliceElements; + uint equation; + uint uintsPerElement; +}; + +kernel void detile_cs( + device const uint* tiled [[buffer(0)]], + device const uint* xTermOrTable [[buffer(1)]], + device const uint* yTerm [[buffer(2)]], + device uint* outLinear [[buffer(3)]], + constant DetileParams& params [[buffer(4)]], + uint3 gid [[thread_position_in_grid]]) +{ + uint elemX = gid.x / params.uintsPerElement; + uint word = gid.x - elemX * params.uintsPerElement; + if (elemX >= params.width || gid.y >= params.height) + { + return; + } + + uint blockIndex = (gid.y / params.blockHeight) * params.blocksPerRow + (elemX / params.blockWidth); + uint off; + if (params.equation != 0u) + { + uint inX = elemX - (elemX / params.blockWidth) * params.blockWidth; + uint inY = gid.y - (gid.y / params.blockHeight) * params.blockHeight; + off = xTermOrTable[inY * params.blockWidth + inX]; + } + else + { + off = xTermOrTable[elemX & params.xMask] ^ yTerm[gid.y & params.yMask]; + } + + uint srcElem = gid.z * params.srcSliceElements + blockIndex * params.blockElements + off; + uint dstElem = gid.z * params.width * params.height + gid.y * params.width + elemX; + outLinear[dstElem * params.uintsPerElement + word] = tiled[srcElem * params.uintsPerElement + word]; +} diff --git a/src/SharpEmu.ShaderCompiler.Vulkan/SpirvFixedShaders.cs b/src/SharpEmu.ShaderCompiler.Vulkan/SpirvFixedShaders.cs index 479a27e3..3b2493c1 100644 --- a/src/SharpEmu.ShaderCompiler.Vulkan/SpirvFixedShaders.cs +++ b/src/SharpEmu.ShaderCompiler.Vulkan/SpirvFixedShaders.cs @@ -260,4 +260,233 @@ public static class SpirvFixedShaders module.AddExecutionMode(main, SpirvExecutionMode.OriginUpperLeft); return module.Build(); } + + /// + /// Compute kernel that deswizzles RDNA2 tiled surfaces at 4 bytes/element into + /// a linear output buffer — one GPU thread per texel, one dispatch-Z layer per + /// array slice. Mirrors GnmTiling.GetDetileParams so it is bit-identical + /// to the CPU fallback for both supported equation families: + /// + /// z = layer; + /// inBlock = equation == BlockTable // modes 1/4/8 + /// ? blockTable[(y % blockHeight) * blockWidth + (x % blockWidth)] + /// : xTerm[x & xMask] ^ yTerm[y & yMask]; // ExactXor 5/9/24/27 + /// src = z * srcSliceElements + /// + (y / blockHeight * blocksPerRow + x / blockWidth) * blockElements + /// + inBlock; + /// out[z * width * height + y * width + x] = tiled[src]; + /// + /// Each array slice is an independently tiled 2D surface; the caller packs the + /// slices contiguously in the tiled buffer (stride srcSliceElements) and + /// the output ends up layer-major, matching a single multi-layer + /// buffer->image copy. For a non-arrayed texture the caller dispatches a + /// single Z layer with srcSliceElements unused (z == 0). + /// + /// The term tables hold ELEMENT offsets. For ExactXor the caller pre-shifts the + /// byte-unit GetDetileParams terms right by log2(bytesPerElement) (exact at 4bpp + /// since the equation's low two byte-offset bits are 0); for BlockTable the + /// GetDetileParams block table is already in element units. Binding 1 carries + /// xTerm (ExactXor) OR blockTable (BlockTable) — the two equations index + /// different-sized buffers, so the kernel branches and evaluates exactly one. + /// + /// width/height are ELEMENT dims (for block-compressed formats a 4x4 block is + /// one element). Each element spans uintsPerElement = bpp/4 words (4bpp -> 1, + /// 8bpp -> 2, 16bpp -> 4); the X dispatch is widened by that factor so each + /// thread copies one word (elemX = gidX / upe, word = gidX % upe). 1/2 bpp are + /// sub-word and stay on the CPU. + /// + /// Descriptor set 0: binding 0 = tiled uint[], 1 = xTerm/blockTable uint[], + /// 2 = yTerm uint[], 3 = out uint[]. Push constants (11 x uint, offset i*4): + /// width, height, blockWidth, blockHeight, blockElements, blocksPerRow, + /// xMask, yMask, srcSliceElements, equation (0 = ExactXor, 1 = BlockTable), + /// uintsPerElement. Local size 8x8x1; dispatch X = ceil(width*upe/8), + /// Y = ceil(height/8), Z = arrayLayers. + /// + public static byte[] CreateDetileCompute() + { + var module = new SpirvModuleBuilder(); + module.AddCapability(SpirvCapability.Shader); + + var voidType = module.TypeVoid(); + var boolType = module.TypeBool(); + var uintType = module.TypeInt(32, signed: false); + var uvec3Type = module.TypeVector(uintType, 3); + + // One shared Block-decorated storage-buffer struct: struct { uint data[]; }. + var runtimeArray = module.TypeRuntimeArray(uintType); + module.AddDecoration(runtimeArray, SpirvDecoration.ArrayStride, 4); + var bufferStruct = module.TypeStruct(runtimeArray); + module.AddDecoration(bufferStruct, SpirvDecoration.Block); + module.AddMemberDecoration(bufferStruct, 0, SpirvDecoration.Offset, 0); + var bufferPtrType = module.TypePointer(SpirvStorageClass.StorageBuffer, bufferStruct); + var uintStoragePtr = module.TypePointer(SpirvStorageClass.StorageBuffer, uintType); + + uint MakeBuffer(uint binding, string name) + { + var variable = module.AddGlobalVariable(bufferPtrType, SpirvStorageClass.StorageBuffer); + module.AddName(variable, name); + module.AddDecoration(variable, SpirvDecoration.DescriptorSet, 0); + module.AddDecoration(variable, SpirvDecoration.Binding, binding); + return variable; + } + + var tiledVar = MakeBuffer(0, "tiled"); + var xTermVar = MakeBuffer(1, "xTerm"); + var yTermVar = MakeBuffer(2, "yTerm"); + var outVar = MakeBuffer(3, "outLinear"); + + // Push constants: struct { uint p0..p10; }, each member at offset i*4. + var pushStruct = module.TypeStruct( + uintType, uintType, uintType, uintType, uintType, uintType, + uintType, uintType, uintType, uintType, uintType); + module.AddDecoration(pushStruct, SpirvDecoration.Block); + for (uint member = 0; member < 11; member++) + { + module.AddMemberDecoration(pushStruct, member, SpirvDecoration.Offset, member * 4); + } + + var pushPtrType = module.TypePointer(SpirvStorageClass.PushConstant, pushStruct); + var pushMemberPtrType = module.TypePointer(SpirvStorageClass.PushConstant, uintType); + var pushVar = module.AddGlobalVariable(pushPtrType, SpirvStorageClass.PushConstant); + module.AddName(pushVar, "pc"); + + var inputUvec3Ptr = module.TypePointer(SpirvStorageClass.Input, uvec3Type); + var gidVar = module.AddGlobalVariable(inputUvec3Ptr, SpirvStorageClass.Input); + module.AddName(gidVar, "gid"); + module.AddDecoration(gidVar, SpirvDecoration.BuiltIn, (uint)SpirvBuiltIn.GlobalInvocationId); + + var uintConst = new uint[11]; + for (uint value = 0; value < 11; value++) + { + uintConst[value] = module.Constant(uintType, value); + } + + var functionType = module.TypeFunction(voidType); + var main = module.BeginFunction(voidType, functionType); + module.AddName(main, "main"); + module.AddLabel(); + + var gid = module.AddInstruction(SpirvOp.Load, uvec3Type, gidVar); + var gidX = module.AddInstruction(SpirvOp.CompositeExtract, uintType, gid, 0); + var y = module.AddInstruction(SpirvOp.CompositeExtract, uintType, gid, 1); + var z = module.AddInstruction(SpirvOp.CompositeExtract, uintType, gid, 2); + + uint PushField(uint index) + { + var pointer = module.AddInstruction( + SpirvOp.AccessChain, pushMemberPtrType, pushVar, uintConst[index]); + return module.AddInstruction(SpirvOp.Load, uintType, pointer); + } + + // width/height are ELEMENT dims (for BC, a 4x4 block is one element). Each + // element spans uintsPerElement 32-bit words (bpp/4: 4bpp->1, 8bpp->2, + // 16bpp->4). The X dispatch is widened by uintsPerElement so each thread + // copies exactly one word: elemX = gidX / upe, wordIndex = gidX % upe. + var width = PushField(0); + var height = PushField(1); + var blockWidth = PushField(2); + var blockHeight = PushField(3); + var blockElements = PushField(4); + var blocksPerRow = PushField(5); + var xMask = PushField(6); + var yMask = PushField(7); + var srcSliceElements = PushField(8); + var equation = PushField(9); + var uintsPerElement = PushField(10); + + var elemX = module.AddInstruction(SpirvOp.UDiv, uintType, gidX, uintsPerElement); + var elemXTimesUpe = module.AddInstruction(SpirvOp.IMul, uintType, elemX, uintsPerElement); + var wordIndex = module.AddInstruction(SpirvOp.ISub, uintType, gidX, elemXTimesUpe); + + var xInRange = module.AddInstruction(SpirvOp.ULessThan, boolType, elemX, width); + var yInRange = module.AddInstruction(SpirvOp.ULessThan, boolType, y, height); + var inRange = module.AddInstruction(SpirvOp.LogicalAnd, boolType, xInRange, yInRange); + + var bodyLabel = module.AllocateId(); + var mergeLabel = module.AllocateId(); + module.AddStatement(SpirvOp.SelectionMerge, mergeLabel, 0); + module.AddStatement(SpirvOp.BranchConditional, inRange, bodyLabel, mergeLabel); + + module.AddLabel(bodyLabel); + + // blockIdx = (y / blockHeight) * blocksPerRow + (elemX / blockWidth) + var yDiv = module.AddInstruction(SpirvOp.UDiv, uintType, y, blockHeight); + var blockRow = module.AddInstruction(SpirvOp.IMul, uintType, yDiv, blocksPerRow); + var xDiv = module.AddInstruction(SpirvOp.UDiv, uintType, elemX, blockWidth); + var blockIdx = module.AddInstruction(SpirvOp.IAdd, uintType, blockRow, xDiv); + + // off (element offset within the block) = equation == BlockTable + // ? blockTable[(y % blockHeight) * blockWidth + (elemX % blockWidth)] + // : xTerm[elemX & xMask] ^ yTerm[y & yMask] + // Binding 1 (xTermVar) doubles as the block table; the two equations index + // different-sized buffers, so exactly one branch executes (no OOB read). + var isBlockTable = module.AddInstruction(SpirvOp.INotEqual, boolType, equation, uintConst[0]); + var xorLabel = module.AllocateId(); + var tableLabel = module.AllocateId(); + var offMergeLabel = module.AllocateId(); + module.AddStatement(SpirvOp.SelectionMerge, offMergeLabel, 0); + module.AddStatement(SpirvOp.BranchConditional, isBlockTable, tableLabel, xorLabel); + + // ExactXor: xTerm[elemX & xMask] ^ yTerm[y & yMask] + module.AddLabel(xorLabel); + var xIdx = module.AddInstruction(SpirvOp.BitwiseAnd, uintType, elemX, xMask); + var xPtr = module.AddInstruction(SpirvOp.AccessChain, uintStoragePtr, xTermVar, uintConst[0], xIdx); + var xTerm = module.AddInstruction(SpirvOp.Load, uintType, xPtr); + var yIdx = module.AddInstruction(SpirvOp.BitwiseAnd, uintType, y, yMask); + var yPtr = module.AddInstruction(SpirvOp.AccessChain, uintStoragePtr, yTermVar, uintConst[0], yIdx); + var yTerm = module.AddInstruction(SpirvOp.Load, uintType, yPtr); + var offXor = module.AddInstruction(SpirvOp.BitwiseXor, uintType, xTerm, yTerm); + module.AddStatement(SpirvOp.Branch, offMergeLabel); + + // BlockTable: blockTable[inY * blockWidth + inX], inX/inY = position in block + module.AddLabel(tableLabel); + var blockXBase = module.AddInstruction(SpirvOp.IMul, uintType, xDiv, blockWidth); + var inX = module.AddInstruction(SpirvOp.ISub, uintType, elemX, blockXBase); + var blockYBase = module.AddInstruction(SpirvOp.IMul, uintType, yDiv, blockHeight); + var inY = module.AddInstruction(SpirvOp.ISub, uintType, y, blockYBase); + var rowInBlock = module.AddInstruction(SpirvOp.IMul, uintType, inY, blockWidth); + var tableIdx = module.AddInstruction(SpirvOp.IAdd, uintType, rowInBlock, inX); + var tablePtr = module.AddInstruction(SpirvOp.AccessChain, uintStoragePtr, xTermVar, uintConst[0], tableIdx); + var offTable = module.AddInstruction(SpirvOp.Load, uintType, tablePtr); + module.AddStatement(SpirvOp.Branch, offMergeLabel); + + module.AddLabel(offMergeLabel); + var off = module.AddInstruction(SpirvOp.Phi, uintType, offXor, xorLabel, offTable, tableLabel); + + // srcElem = z * srcSliceElements + blockIdx * blockElements + off (in elements) + // srcWord = srcElem * uintsPerElement + wordIndex + var srcSliceBase = module.AddInstruction(SpirvOp.IMul, uintType, z, srcSliceElements); + var blockBase = module.AddInstruction(SpirvOp.IMul, uintType, blockIdx, blockElements); + var srcInSlice = module.AddInstruction(SpirvOp.IAdd, uintType, blockBase, off); + var srcElem = module.AddInstruction(SpirvOp.IAdd, uintType, srcSliceBase, srcInSlice); + var srcElemWords = module.AddInstruction(SpirvOp.IMul, uintType, srcElem, uintsPerElement); + var src = module.AddInstruction(SpirvOp.IAdd, uintType, srcElemWords, wordIndex); + var srcPtr = module.AddInstruction(SpirvOp.AccessChain, uintStoragePtr, tiledVar, uintConst[0], src); + var word = module.AddInstruction(SpirvOp.Load, uintType, srcPtr); + + // dstElem = z * width * height + y * width + elemX (in elements) + // dstWord = dstElem * uintsPerElement + wordIndex + var sliceElements = module.AddInstruction(SpirvOp.IMul, uintType, width, height); + var dstSliceBase = module.AddInstruction(SpirvOp.IMul, uintType, z, sliceElements); + var rowBase = module.AddInstruction(SpirvOp.IMul, uintType, y, width); + var dstRow = module.AddInstruction(SpirvOp.IAdd, uintType, rowBase, elemX); + var dstElem = module.AddInstruction(SpirvOp.IAdd, uintType, dstSliceBase, dstRow); + var dstElemWords = module.AddInstruction(SpirvOp.IMul, uintType, dstElem, uintsPerElement); + var dstIdx = module.AddInstruction(SpirvOp.IAdd, uintType, dstElemWords, wordIndex); + var dstPtr = module.AddInstruction(SpirvOp.AccessChain, uintStoragePtr, outVar, uintConst[0], dstIdx); + module.AddStatement(SpirvOp.Store, dstPtr, word); + + module.AddStatement(SpirvOp.Branch, mergeLabel); + module.AddLabel(mergeLabel); + module.AddStatement(SpirvOp.Return); + module.EndFunction(); + + module.AddExecutionMode(main, SpirvExecutionMode.LocalSize, 8, 8, 1); + module.AddEntryPoint( + SpirvExecutionModel.GLCompute, + main, + "main", + [gidVar, tiledVar, xTermVar, yTermVar, outVar, pushVar]); + return module.Build(); + } } diff --git a/tests/SharpEmu.Libs.Tests/Agc/DetileComputeSpirvTests.cs b/tests/SharpEmu.Libs.Tests/Agc/DetileComputeSpirvTests.cs new file mode 100644 index 00000000..76e57e2a --- /dev/null +++ b/tests/SharpEmu.Libs.Tests/Agc/DetileComputeSpirvTests.cs @@ -0,0 +1,102 @@ +// Copyright (C) 2026 SharpEmu Emulator Project +// SPDX-License-Identifier: GPL-2.0-or-later + +using System.Buffers.Binary; +using SharpEmu.ShaderCompiler.Vulkan; +using Xunit; + +namespace SharpEmu.Libs.Tests.Agc; + +// Structural validation of the GPU detile compute kernel. This cannot run the +// shader without a Vulkan device, but it pins the SPIR-V is well-formed: the +// header is correct, every instruction's word count sums to exactly the module +// length (the classic hand-emit bug), and the compute-specific pieces are +// present (a GLCompute entry point, a LocalSize execution mode, and a runtime +// array for the storage buffers). Full pixel correctness is verified on a GPU. +public sealed class DetileComputeSpirvTests +{ + private const uint SpirvMagic = 0x07230203; + private const uint SpirvVersion15 = 0x00010500; + + private const ushort OpEntryPoint = 15; + private const ushort OpExecutionMode = 16; + private const ushort OpTypeRuntimeArray = 29; + private const ushort OpFunction = 54; + private const ushort OpFunctionEnd = 56; + + private const uint ExecutionModelGLCompute = 5; + private const uint ExecutionModeLocalSize = 17; + + [Fact] + public void CreateDetileCompute_EmitsWellFormedComputeModule() + { + var spirv = SpirvFixedShaders.CreateDetileCompute(); + + Assert.True(spirv.Length % sizeof(uint) == 0, "SPIR-V must be a whole number of words."); + var words = new uint[spirv.Length / sizeof(uint)]; + for (var i = 0; i < words.Length; i++) + { + words[i] = BinaryPrimitives.ReadUInt32LittleEndian(spirv.AsSpan(i * sizeof(uint))); + } + + Assert.True(words.Length > 5, "Module must have a header plus instructions."); + Assert.Equal(SpirvMagic, words[0]); + Assert.Equal(SpirvVersion15, words[1]); + + var bound = words[3]; + Assert.True(bound > 1, "Id bound must be set."); + + var sawComputeEntry = false; + var sawLocalSize = false; + var sawRuntimeArray = false; + var functionCount = 0; + var functionEndCount = 0; + + var offset = 5; + while (offset < words.Length) + { + var word = words[offset]; + var wordCount = (int)(word >> 16); + var opcode = (ushort)(word & 0xFFFF); + + Assert.True(wordCount >= 1, $"Instruction at {offset} has a zero word count."); + Assert.True( + offset + wordCount <= words.Length, + $"Instruction at {offset} (op {opcode}, wc {wordCount}) overruns the module."); + + switch (opcode) + { + case OpEntryPoint when words[offset + 1] == ExecutionModelGLCompute: + sawComputeEntry = true; + break; + case OpExecutionMode + when wordCount >= 6 && + words[offset + 2] == ExecutionModeLocalSize && + words[offset + 3] == 8 && + words[offset + 4] == 8 && + words[offset + 5] == 1: + sawLocalSize = true; + break; + case OpTypeRuntimeArray: + sawRuntimeArray = true; + break; + case OpFunction: + functionCount++; + break; + case OpFunctionEnd: + functionEndCount++; + break; + } + + offset += wordCount; + } + + // Word counts must tile the module exactly — a wrong length lands here. + Assert.Equal(words.Length, offset); + Assert.True(sawComputeEntry, "Missing a GLCompute OpEntryPoint."); + Assert.True(sawLocalSize, "Missing an 8x8x1 LocalSize execution mode."); + Assert.True(sawRuntimeArray, "Missing a runtime array (storage buffers)."); + Assert.Equal(1, functionCount); + Assert.Equal(1, functionEndCount); + } +} diff --git a/tests/SharpEmu.Libs.Tests/Agc/GnmTilingDetileTests.cs b/tests/SharpEmu.Libs.Tests/Agc/GnmTilingDetileTests.cs index b517f53c..d33ef96d 100644 --- a/tests/SharpEmu.Libs.Tests/Agc/GnmTilingDetileTests.cs +++ b/tests/SharpEmu.Libs.Tests/Agc/GnmTilingDetileTests.cs @@ -83,4 +83,152 @@ public sealed class GnmTilingDetileTests Assert.Equal((ushort)i, value); } } + + // GetDetileParams must reproduce TryDetile bit-for-bit: the CPU fallback and + // the GPU compute kernel both consume these params, so a detile driven purely + // by DetileParams (the shared addressing formula the kernel runs) must equal + // the shipped CPU detile for every supported mode/bpp. + [Theory] + [InlineData(27u, 2, 384, 200)] // 64 KiB RB+ R_X (exact-XOR) + [InlineData(27u, 4, 256, 256)] // 64 KiB RB+ R_X (exact-XOR) + [InlineData(9u, 4, 300, 300)] // 64 KiB standard (exact-XOR) + [InlineData(24u, 4, 128, 256)] // 64 KiB RB+ Z_X (exact-XOR) + [InlineData(5u, 4, 200, 120)] // 4 KiB standard (exact-XOR) + [InlineData(8u, 4, 128, 128)] // 64 KiB Z (block-table path) + [InlineData(1u, 4, 64, 64)] // 256 B standard (block-table path) + public void GetDetileParams_ReproducesTryDetile(uint mode, int bpp, int w, int h) + { + var p = GnmTiling.GetDetileParams(mode, bpp, w, h); + Assert.True(p.IsSupported); + + // Whole-block tiled buffer (block addressing overshoots the linear extent), + // filled with a deterministic non-trivial pattern. + var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight; + var tiled = new byte[(long)p.BlocksPerRow * blocksHigh * p.BlockBytes]; + for (var i = 0; i < tiled.Length; i++) + { + tiled[i] = (byte)((i * 31 + 7) & 0xFF); + } + + var expected = new byte[w * h * bpp]; + Assert.True(GnmTiling.TryDetile(tiled, expected, mode, w, h, bpp)); + + var actual = DetileViaParams(tiled, p, w, h, bpp); + Assert.Equal(expected, actual); + } + + // The production GnmTiling.DetileWithParams (the active CPU fallback used by + // the Metal path under default-on GPU detile) must equal TryDetile for every + // supported mode/bpp — same DetileParams addressing, no re-derived swizzle. + [Theory] + [InlineData(27u, 2, 384, 200)] + [InlineData(27u, 4, 256, 256)] + [InlineData(9u, 4, 300, 300)] + [InlineData(24u, 4, 128, 256)] + [InlineData(5u, 4, 200, 120)] + [InlineData(8u, 4, 128, 128)] + [InlineData(1u, 4, 64, 64)] + [InlineData(27u, 8, 256, 256)] // 8bpp (GPU: 2 words/element) + [InlineData(27u, 16, 128, 128)] // 16bpp (GPU: 4 words/element) + [InlineData(9u, 8, 128, 96)] + [InlineData(8u, 16, 64, 64)] // block-table, 16bpp + public void DetileWithParams_MatchesTryDetile(uint mode, int bpp, int w, int h) + { + var p = GnmTiling.GetDetileParams(mode, bpp, w, h); + Assert.True(p.IsSupported); + + var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight; + var tiled = new byte[(long)p.BlocksPerRow * blocksHigh * p.BlockBytes]; + for (var i = 0; i < tiled.Length; i++) + { + tiled[i] = (byte)((i * 31 + 7) & 0xFF); + } + + var expected = new byte[w * h * bpp]; + Assert.True(GnmTiling.TryDetile(tiled, expected, mode, w, h, bpp)); + + var actual = new byte[w * h * bpp]; + Assert.True(GnmTiling.DetileWithParams(p, tiled, actual)); + Assert.Equal(expected, actual); + } + + // Array textures are packed as contiguous tiled slices and detiled one slice + // per layer (dispatch-Z on the GPU; a per-layer loop in the CPU fallbacks) + // into a layer-major linear buffer. This pins that packing: each slice must + // deswizzle into its own region and match a per-slice TryDetile, and a + // per-layer-distinct pattern catches any slice cross-talk. + [Theory] + [InlineData(27u, 4, 256, 256, 3)] + [InlineData(9u, 4, 128, 96, 2)] + [InlineData(24u, 4, 64, 128, 4)] + public void DetileWithParams_MultiLayer_MatchesPerSliceTryDetile(uint mode, int bpp, int w, int h, int layers) + { + var p = GnmTiling.GetDetileParams(mode, bpp, w, h); + Assert.True(p.IsSupported); + + var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight; + var sliceTiledBytes = (int)((long)p.BlocksPerRow * blocksHigh * p.BlockBytes); + var sliceLinearBytes = w * h * bpp; + + var tiled = new byte[sliceTiledBytes * layers]; + for (var layer = 0; layer < layers; layer++) + { + for (var i = 0; i < sliceTiledBytes; i++) + { + tiled[layer * sliceTiledBytes + i] = (byte)((i * 31 + 7 + layer * 101) & 0xFF); + } + } + + // Expected: each slice detiled independently via the shipped CPU detile. + var expected = new byte[sliceLinearBytes * layers]; + for (var layer = 0; layer < layers; layer++) + { + Assert.True(GnmTiling.TryDetile( + tiled.AsSpan(layer * sliceTiledBytes, sliceTiledBytes), + expected.AsSpan(layer * sliceLinearBytes, sliceLinearBytes), + mode, w, h, bpp)); + } + + // Actual: the layer-major loop the Vulkan/Metal CPU fallbacks run. + var actual = new byte[sliceLinearBytes * layers]; + for (var layer = 0; layer < layers; layer++) + { + Assert.True(GnmTiling.DetileWithParams( + p, + tiled.AsSpan(layer * sliceTiledBytes, sliceTiledBytes), + actual.AsSpan(layer * sliceLinearBytes, sliceLinearBytes))); + } + + Assert.Equal(expected, actual); + } + + // Reference detile driven entirely by DetileParams — the single shared + // addressing formula the Vulkan/Metal compute kernel will run per texel. + private static byte[] DetileViaParams(byte[] tiled, DetileParams p, int w, int h, int bpp) + { + var linear = new byte[w * h * bpp]; + for (var y = 0; y < h; y++) + { + for (var x = 0; x < w; x++) + { + var blockX = x / p.BlockWidth; + var blockY = y / p.BlockHeight; + var inX = x % p.BlockWidth; + var inY = y % p.BlockHeight; + var inBlockByte = p.Equation == DetileEquation.ExactXor + ? p.XByteTerm[x & p.XMask] ^ p.YByteTerm[y & p.YMask] + : p.BlockTable[inY * p.BlockWidth + inX] * p.BytesPerElement; + var srcByte = ((long)blockY * p.BlocksPerRow + blockX) * p.BlockBytes + inBlockByte; + var dstByte = ((long)y * w + x) * bpp; + if (srcByte < 0 || srcByte + bpp > tiled.Length) + { + continue; + } + + Array.Copy(tiled, srcByte, linear, dstByte, bpp); + } + } + + return linear; + } }