Compare commits

..

4 Commits

Author SHA1 Message Date
ParantezTech 77fe07efc0 [Bink2] rework bridge to use FFmpeg's native Bink2 decoder instead of a C bridge 2026-07-23 02:52:24 +03:00
ParantezTech 5f37dd85e0 [CMake] update commit 2026-07-22 20:28:36 +03:00
ParantezTech 38621e7be9 [CMake] update commit 2026-07-22 19:46:19 +03:00
ParantezTech c9c0793059 [VideoOut] Add Bink2 support via FFMPEG bridge 2026-07-22 18:31:29 +03:00
53 changed files with 551 additions and 7817 deletions
Binary file not shown.

Before

Width:  |  Height:  |  Size: 345 KiB

After

Width:  |  Height:  |  Size: 190 KiB

Binary file not shown.

Before

Width:  |  Height:  |  Size: 229 KiB

After

Width:  |  Height:  |  Size: 82 KiB

Binary file not shown.

Before

Width:  |  Height:  |  Size: 227 KiB

After

Width:  |  Height:  |  Size: 101 KiB

Binary file not shown.

Before

Width:  |  Height:  |  Size: 104 KiB

After

Width:  |  Height:  |  Size: 86 KiB

-19
View File
@@ -26,25 +26,6 @@ Before opening a pull request, please keep the following in mind:
If you're unsure about a design decision, feel free to open a discussion or draft PR first.
## Pull Request Expectations
Pull requests should provide real, observable emulator behavior rather than only suppressing errors or unresolved imports.
Changes that only return success, zero, or fabricated handles without implementing the expected state, output, or side effects will generally not be accepted. Functions that create resources, write output structures, register callbacks, or expose runtime state should model the behavior required by the guest.
When applicable, PRs should include:
- The affected game or application.
- Relevant logs or failing imports.
- Behavior before and after the change.
- Real game testing and known limitations.
Avoid submitting large collections of speculative NIDs or unrelated exports. Keep each PR focused on one problem or a closely related set of changes.
Large architectural changes should be discussed with the maintainers before implementation. Contributors are encouraged to ask first when they are uncertain whether a proposed direction fits the project.
Opening a PR does not guarantee that it will be merged. Maintainers evaluate changes based on correctness, evidence, testing, scope, maintenance cost, and the long-term direction of the project.
## AI-Assisted Contributions
AI-assisted development is welcome and may be used for research, reverse engineering, code generation, or documentation.
+1 -1
View File
@@ -9,7 +9,7 @@ SPDX-License-Identifier: GPL-2.0-or-later
<ImplicitUsings>enable</ImplicitUsings>
<Nullable>enable</Nullable>
<GenerateDocumentationFile>true</GenerateDocumentationFile>
<SharpEmuVersion>0.0.2-beta.5</SharpEmuVersion>
<SharpEmuVersion>0.0.2-beta.4</SharpEmuVersion>
<Version>$(SharpEmuVersion)</Version>
<RepoRoot>$([MSBuild]::NormalizeDirectory('$(MSBuildThisFileDirectory)'))</RepoRoot>
-20
View File
@@ -25,14 +25,6 @@ SPDX-License-Identifier: GPL-2.0-or-later
---
<p align="center">
<a href="#support">
<img src="https://img.shields.io/badge/Support-GitHub%20Sponsors%20%26%20Crypto-EA4AAA?style=for-the-badge&logo=githubsponsors&logoColor=white" alt="Support SharpEmu">
</a>
</p>
---
> [!NOTE]
> SharpEmu supports Windows x64, Linux x64, and macOS x64. Apple Silicon Macs
> can run the macOS x64 build through Rosetta 2, and Windows on ARM devices
@@ -144,18 +136,6 @@ Provided valuable references for filesystem handling and low-level C# implementa
- [**GPL-2.0 license**](https://github.com/sharpemu/sharpemu/blob/main/LICENSE)
## Support
Support SharpEmu via GitHub Sponsors or cryptocurrency. Every contribution helps fund ongoing development and long-term maintenance. GitHub Sponsors is the preferred way to support the project, but cryptocurrency donations are also appreciated.
### ETH/USDT
`0xF315F5d986c790bB3A58DbE60F1B2760997dEd82`
### BTC
`bc1qmr9k8899njys5ny63xsues4jgmkk96erslrkmv`
## Contributing
Before opening an issue or pull request, please read our contribution guidelines:
-1
View File
@@ -21,7 +21,6 @@ SPDX-License-Identifier: GPL-2.0-or-later
<Folder Name="/tests/">
<Project Path="tests/SharpEmu.Libs.Tests/SharpEmu.Libs.Tests.csproj" />
<Project Path="tests/SharpEmu.ShaderCompiler.Metal.Tests/SharpEmu.ShaderCompiler.Metal.Tests.csproj" />
<Project Path="tests/SharpEmu.ShaderCompiler.Tests/SharpEmu.ShaderCompiler.Tests.csproj" />
<Project Path="tests/SharpEmu.SourceGenerators.Tests/SharpEmu.SourceGenerators.Tests.csproj" />
</Folder>
</Solution>
-1
View File
@@ -153133,7 +153133,6 @@ scePsmlMfsrGetContextBufferRequirement800M3_2
scePsmlMfsrGetDispatchMfsrPacket1000
scePsmlMfsrGetDispatchMfsrPacket1100
scePsmlMfsrGetDispatchMfsrPacketSizeInDwords
scePsmlMfsrGetDispatchMfsrPacket900
scePsmlMfsrGetMipmapBias
scePsmlMfsrGetSharedResourcesInitRequirement
scePsmlMfsrInit
+6 -1
View File
@@ -573,7 +573,12 @@ internal static partial class Program
return false;
}
string[] childArgs = [MitigatedChildFlag, .. args];
var childArgs = new string[args.Length + 1];
childArgs[0] = MitigatedChildFlag;
for (var i = 0; i < args.Length; i++)
{
childArgs[i + 1] = args[i];
}
var commandLine = BuildCommandLine(processPath, childArgs);
var startupInfoEx = new STARTUPINFOEX();
@@ -19,9 +19,6 @@ public sealed partial class DirectExecutionBackend
private static int _lazyCommitTraceCount;
private static int _guestAllocatorHoleRecoveries;
private static int _auxiliaryThreadExecuteFaultRecoveries;
private static int _auxiliaryThreadExecuteFaultSkips;
private nint _workerAbortStack;
private const uint WorkerAbortStackSize = 0x10000u;
private unsafe void SetupExceptionHandler()
{
@@ -438,91 +435,18 @@ public sealed partial class DirectExecutionBackend
void* contextRecord,
ulong rip)
{
if (exceptionRecord->ExceptionCode != 3221225477u)
if (exceptionRecord->ExceptionCode != 3221225477u ||
rip >= 0x0000000800000000UL ||
_activeGuestThreadState is not { Name: "tbb_thead" } activeThread)
{
return false;
}
// Prefer ThreadStatic active state; fall back to host-thread name when
// concurrent TBB AVs race logging (tLT61: recover skipped, then Fatal).
GuestThreadState? activeThread = _activeGuestThreadState;
if (activeThread is null || activeThread.Name != "tbb_thead")
{
var hostName = Thread.CurrentThread.Name;
if (hostName is null ||
!hostName.StartsWith("SharpEmu-tbb_thead", StringComparison.Ordinal))
{
return false;
}
activeThread = FindGuestThreadStateByHostThreadId(unchecked((int)GetCurrentThreadId()));
if (activeThread is null || activeThread.Name != "tbb_thead")
{
var skip = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultSkips);
if (skip <= 8 || skip % 64 == 0)
{
Console.Error.WriteLine(
$"[LOADER][WARN] tbb_recover skip #{skip}: rip=0x{rip:X16} " +
$"host='{hostName}' active={(activeThread?.Name ?? "null")}");
Console.Error.Flush();
}
return false;
}
}
var hostExit = ActiveEntryReturnSentinelRip;
if (hostExit < 0x10000)
{
hostExit = unchecked((ulong)_guestReturnStub);
}
// Prefer worker-abort (SetEvent + ExitThread) over host_exit→RunEpilogue:
// the latter FailFasts the process after TBB recover (tLT28/30 silent die).
// Do NOT abandon mutexes here — managed HLE from inside VEH can re-enter
// and Fatal (tLT73). NativeGuestExecutor.Run abandons after detecting abort.
var abortRip = unchecked((ulong)_workerAbortStub);
if (abortRip >= 0x10000)
{
// Do NOT SetEvent from managed VEH: that wakes the renter which may
// TerminateThread while this thread is still inside VEH return
// (tLTA2: recover logged, no respawning, process die). Abort stub
// SetEvent's only after CONTINUE_EXECUTION resumes at park.
// Prefer the entry-stub-saved host RSP (real CreateThread stack).
// Do not treat mid-range host stacks as guest — Astro worker stacks
// often sit in 0x02xxxxxx_xxxx and were wrongly replaced with a
// shared VirtualAlloc abort stack (concurrent TBB AV → die).
var hostRspSlot = TlsGetValue(_hostRspSlotTlsIndex);
ulong hostRsp = 0;
if (hostRspSlot != 0)
{
hostRsp = *(ulong*)hostRspSlot;
}
if (hostRsp < 0x10000)
{
hostRsp = EnsureWorkerAbortStackRsp();
}
if (hostRsp >= 0x10000)
{
WriteCtxU64(contextRecord, 152, hostRsp & ~0xFUL);
}
WriteCtxU64(contextRecord, 120, 0);
WriteCtxU64(contextRecord, 248, abortRip);
var recovery = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultRecoveries);
Console.Error.WriteLine(
$"[LOADER][WARN] Recovered auxiliary TBB execute fault #{recovery}: " +
$"thread=0x{activeThread.ThreadHandle:X16} target=0x{rip:X16} " +
$"host_rsp=0x{hostRsp:X16} -> worker_abort=0x{abortRip:X16}");
Console.Error.WriteLine(
"[LOADER][INFO] tbb_recover: parking native worker (SetEvent+park); " +
"renter will TerminateThread+respawn — avoids ExitThread after VEH");
Console.Error.Flush();
return true;
}
if (hostExit < 0x10000)
{
Console.Error.WriteLine(
@@ -534,57 +458,13 @@ public sealed partial class DirectExecutionBackend
_ = TryPatchActiveGuestReturnSlot(hostExit);
WriteCtxU64(contextRecord, 120, 0);
WriteCtxU64(contextRecord, 248, hostExit);
var recoveryFallback = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultRecoveries);
var recovery = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultRecoveries);
Console.Error.WriteLine(
$"[LOADER][WARN] Recovered auxiliary TBB execute fault #{recoveryFallback}: " +
$"[LOADER][WARN] Recovered auxiliary TBB execute fault #{recovery}: " +
$"thread=0x{activeThread.ThreadHandle:X16} target=0x{rip:X16} -> host_exit=0x{hostExit:X16}");
Console.Error.WriteLine(
"[LOADER][INFO] tbb_recover: resumed at host_exit (abort stub unavailable); " +
"subsequent FastFail/CLR must not re-enter managed VEH " +
"(live trampoline pre-filters 0xC0000409 / 0xE0434352)");
Console.Error.Flush();
return true;
}
private GuestThreadState? FindGuestThreadStateByHostThreadId(int hostThreadId)
{
if (hostThreadId == 0)
{
return null;
}
try
{
foreach (var thread in SnapshotGuestThreads())
{
if (Volatile.Read(ref thread.HostThreadId) == hostThreadId)
{
return thread;
}
}
}
catch
{
}
return null;
}
private unsafe ulong EnsureWorkerAbortStackRsp()
{
if (_workerAbortStack == 0)
{
_workerAbortStack = (nint)VirtualAlloc(null, WorkerAbortStackSize, 12288u, 4u);
if (_workerAbortStack == 0)
{
return 0;
}
}
// Grow-down stack: hand out near the top with alignment headroom.
return (ulong)(_workerAbortStack + (nint)WorkerAbortStackSize - 0x100) & ~0xFUL;
}
private unsafe bool TryRecoverGuestInt41(uint exceptionCode, void* contextRecord, ulong rip)
{
if (!_ignoreGuestInt41 || exceptionCode != 3221225477u || rip < 0x10000)
@@ -1449,9 +1449,6 @@ public sealed partial class DirectExecutionBackend
var expectedSemaphoreTrywaitAgain =
string.Equals(nid, "H2a+IN9TP0E", StringComparison.Ordinal) &&
result == OrbisGen2Result.ORBIS_GEN2_ERROR_TRY_AGAIN;
var expectedPollSemaBusy =
string.Equals(nid, "12wOHk8ywb0", StringComparison.Ordinal) &&
result == OrbisGen2Result.ORBIS_GEN2_ERROR_BUSY;
var expectedNetAcceptWouldBlock =
string.Equals(nid, "PIWqhn9oSxc", StringComparison.Ordinal) &&
resultValue == unchecked((int)0x80410123);
@@ -1466,7 +1463,6 @@ public sealed partial class DirectExecutionBackend
!expectedEqueueTimeout &&
!expectedMutexTrylockBusy &&
!expectedSemaphoreTrywaitAgain &&
!expectedPollSemaBusy &&
!expectedNetAcceptWouldBlock &&
!expectedUserServiceNoEvent &&
!expectedPrivacyInvalidParameter)
@@ -29,29 +29,9 @@ public sealed partial class DirectExecutionBackend
private static readonly bool NativeGuestWorkersDisabled =
string.Equals(Environment.GetEnvironmentVariable("SHARPEMU_DISABLE_NATIVE_GUEST_WORKERS"), "1", StringComparison.Ordinal);
// Cap concurrent native-worker Runs. Astro's tbb_thead burst overlaps many
// UnmanagedCallersOnly prologues; a large prewarm + unbounded concurrency
// FailFasts (0xC0000409) mid-storm with no VEH breadcrumb. Pool size and
// in-flight Runs are separate knobs.
private static readonly int NativeWorkerMaxConcurrent = ReadNativeWorkerMaxConcurrent();
private static int ReadNativeWorkerMaxConcurrent()
{
if (int.TryParse(
Environment.GetEnvironmentVariable("SHARPEMU_NATIVE_WORKER_MAX_CONCURRENT"),
out var parsed) &&
parsed > 0)
{
return Math.Clamp(parsed, 1, 64);
}
return 2;
}
private readonly object _nativeWorkerGate = new();
private readonly List<NativeGuestExecutor> _allNativeWorkers = new();
private readonly Stack<NativeGuestExecutor> _idleNativeWorkers = new();
private readonly SemaphoreSlim _nativeWorkerRunLimiter = new(NativeWorkerMaxConcurrent);
private bool _nativeWorkersDisposed;
private int _nativeWorkerCreationFailedLogged;
@@ -69,9 +49,6 @@ public sealed partial class DirectExecutionBackend
[DllImport("kernel32.dll", SetLastError = true)]
private static extern uint WaitForSingleObject(nint hHandle, uint dwMilliseconds);
[DllImport("kernel32.dll", SetLastError = true)]
private static extern bool TerminateThread(nint hThread, uint dwExitCode);
// Runs an emitted guest entry stub. Preferred path is a pooled native worker
// thread; falls back to the historical inline calli (guest frames above this
// thread's managed frames) when workers are disabled or unavailable.
@@ -79,148 +56,40 @@ public sealed partial class DirectExecutionBackend
// Callers set the Active* thread-statics before emitting the stub and read the
// yield/forced-exit flags right after this returns, so the worker outcome is
// copied back into this thread's statics before returning.
private unsafe int RunGuestEntryStub(void* entryStub, ulong hostRspSlot, bool requireNativeWorker = false)
private unsafe int RunGuestEntryStub(void* entryStub, ulong hostRspSlot)
{
// Limit in-flight native Runs before renting so the idle pool is not
// drained by threads blocked on the concurrency gate.
_nativeWorkerRunLimiter.Wait();
NativeGuestExecutor? worker = null;
var worker = RentNativeGuestExecutor();
if (worker is null)
{
TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot);
return CallNativeEntry(entryStub);
}
try
{
// Astro can spawn a burst of tbb_thead while workers are still in
// TerminateThread+respawn. Wait for a native worker — never fall back
// to managed inline (FailFast) and never throw (uncaught throw mid-
// storm was a silent process die).
var maxAttempts = requireNativeWorker ? 500 : 48;
for (var attempt = 0; attempt < maxAttempts; attempt++)
{
worker = RentNativeGuestExecutor();
if (worker is not null)
{
break;
}
if (!requireNativeWorker)
{
break;
}
Thread.Sleep(attempt < 32 ? 1 : 4);
}
if (worker is null)
{
if (requireNativeWorker)
{
var n = Interlocked.Increment(ref _tbbNativeWorkerRefuseCount);
if (n <= 8 || n % 32 == 0)
{
Console.Error.WriteLine(
$"[LOADER][ERROR] tbb_native_worker unavailable #{n} after {maxAttempts} attempts; " +
"skipping run (no managed inline, no throw)");
Console.Error.Flush();
}
_activeGuestThreadYieldRequested = true;
_activeGuestThreadYieldReason = "tbb_native_worker_unavailable";
_activeForcedGuestExit = true;
return unchecked((int)0x80020012);
}
TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot);
return CallNativeEntry(entryStub);
}
try
{
var state = _activeGuestThreadState;
if (state is { Name: "tbb_thead" })
{
var n = Interlocked.Increment(ref _tbbNativeRunEnterCount);
if (n <= 12 || n % 64 == 0)
{
Console.Error.WriteLine(
$"[LOADER][INFO] tbb_run_enter #{n} native_tid_pending handle=0x{state.ThreadHandle:X16} " +
$"max_concurrent={NativeWorkerMaxConcurrent}");
Console.Error.Flush();
}
}
var nativeReturn = worker.Run(
_activeCpuContext!,
state,
GuestThreadExecution.CurrentGuestThreadHandle,
_activeEntryReturnSentinelRip,
_activeGuestReturnSlotAddress,
(nint)hostRspSlot,
(nint)entryStub,
state?.AffinityMask ?? 0,
out var yieldRequested,
out var yieldReason,
out var forcedExit);
_activeGuestThreadYieldRequested = yieldRequested;
_activeGuestThreadYieldReason = yieldReason;
_activeForcedGuestExit = forcedExit;
return nativeReturn;
}
finally
{
ReturnNativeGuestExecutor(worker);
}
var state = _activeGuestThreadState;
var nativeReturn = worker.Run(
_activeCpuContext!,
state,
GuestThreadExecution.CurrentGuestThreadHandle,
_activeEntryReturnSentinelRip,
_activeGuestReturnSlotAddress,
(nint)hostRspSlot,
(nint)entryStub,
state?.AffinityMask ?? 0,
out var yieldRequested,
out var yieldReason,
out var forcedExit);
_activeGuestThreadYieldRequested = yieldRequested;
_activeGuestThreadYieldReason = yieldReason;
_activeForcedGuestExit = forcedExit;
return nativeReturn;
}
finally
{
_nativeWorkerRunLimiter.Release();
ReturnNativeGuestExecutor(worker);
}
}
private static int _tbbNativeRunEnterCount;
private static int _tbbNativeWorkerRefuseCount;
internal static int _tbbWorkerPrologueFaultCount;
private void PrewarmNativeGuestWorkers(int count)
{
if (!OperatingSystem.IsWindows() || NativeGuestWorkersDisabled || count <= 0)
{
return;
}
var warmed = new List<NativeGuestExecutor>(count);
for (var i = 0; i < count; i++)
{
var worker = NativeGuestExecutor.TryCreate(this);
if (worker is null)
{
break;
}
warmed.Add(worker);
}
lock (_nativeWorkerGate)
{
if (_nativeWorkersDisposed)
{
foreach (var worker in warmed)
{
worker.Dispose();
}
return;
}
foreach (var worker in warmed)
{
_allNativeWorkers.Add(worker);
_idleNativeWorkers.Push(worker);
}
}
Console.Error.WriteLine(
$"[LOADER][INFO] Native guest workers prewarmed: {warmed.Count}/{count} " +
$"max_concurrent={NativeWorkerMaxConcurrent}");
Console.Error.Flush();
}
private NativeGuestExecutor? RentNativeGuestExecutor()
{
// NativeGuestExecutor emits a Win32 wait loop and creates it with
@@ -531,22 +400,6 @@ public sealed partial class DirectExecutionBackend
return false;
}
FlushInstructionCache(GetCurrentProcess(), _loopStub, LoopStubSize);
return StartWorkerThread();
}
private bool RestartWorkerThread()
{
if (_loopStub == null || _controlBlock == null)
{
return false;
}
*(int*)_controlBlock = 0;
return StartWorkerThread();
}
private bool StartWorkerThread()
{
_threadHandle = CreateThread(
0,
WorkerStackReservation,
@@ -592,49 +445,6 @@ public sealed partial class DirectExecutionBackend
_runForcedExit = false;
SignalWorkAvailable();
WaitWorkCompleted();
// Normal path: RunEpilogue/ExitRun clears _entered before SetEvent(done).
// TBB abort stub SetEvent's without ExitRun — _entered stays true.
if (_entered)
{
var waitRc = WaitForSingleObject(_threadHandle, 500u);
Console.Error.WriteLine(
$"[LOADER][WARN] Native guest worker tid={_nativeThreadId} aborted during run; " +
$"wait_rc=0x{waitRc:X8} respawning");
Console.Error.Flush();
if (_runState is { } abortedState)
{
_ = GuestThreadExecution.NotifyGuestThreadAbandoned(
abortedState.ThreadHandle,
"tbb_worker_abort");
Volatile.Write(ref abortedState.HostThreadId, _prevHostThreadId);
}
_entered = false;
if (_threadHandle != 0)
{
// Abort stub parks (no ExitThread). Force-kill the parked OS
// thread so we can recreate the loop without process teardown.
if (waitRc != 0u)
{
_ = TerminateThread(_threadHandle, unchecked((uint)(-1)));
_ = WaitForSingleObject(_threadHandle, 1000u);
}
CloseHandle(_threadHandle);
_threadHandle = 0;
_nativeThreadId = 0;
}
if (!RestartWorkerThread())
{
_runPrologueFailed = true;
}
else
{
_runPrologueFailed = false;
_runForcedExit = true;
_runNativeResult = 0;
}
}
_runContext = null;
_runState = null;
yieldRequested = _runYieldRequested;
@@ -642,22 +452,7 @@ public sealed partial class DirectExecutionBackend
forcedExit = _runForcedExit;
if (_runPrologueFailed)
{
// Never throw out of the native-worker rent path: an uncaught
// exception mid-TBB storm kills the process with no FailFast
// breadcrumb.
var n = Interlocked.Increment(ref _tbbWorkerPrologueFaultCount);
if (n <= 8 || n % 32 == 0)
{
Console.Error.WriteLine(
$"[LOADER][ERROR] tbb_worker prologue fault #{n}; soft-fail run " +
$"(tid={_nativeThreadId})");
Console.Error.Flush();
}
yieldRequested = true;
yieldReason = "tbb_worker_prologue_fault";
forcedExit = true;
return unchecked((int)0x80020012);
throw new InvalidOperationException("Native guest worker failed to bind the run ambient (prologue fault)");
}
return _runNativeResult;
}
@@ -751,18 +546,6 @@ public sealed partial class DirectExecutionBackend
_activeGuestThreadState = _runState;
backend.BindTlsBase(_runContext!);
TlsSetValue(backend._hostRspSlotTlsIndex, _runHostRspSlot);
if (backend._workerDoneEventTlsIndex != uint.MaxValue)
{
nint doneHandle = OperatingSystem.IsWindows()
? _workCompleted!.SafeWaitHandle.DangerousGetHandle()
: _doneSemaphore;
TlsSetValue(backend._workerDoneEventTlsIndex, doneHandle);
}
if (backend._tbbAbortEligibleTlsIndex != uint.MaxValue)
{
nint eligible = _runState is { Name: "tbb_thead" } ? 1 : 0;
TlsSetValue(backend._tbbAbortEligibleTlsIndex, eligible);
}
if (_runState is { } state)
{
_prevHostThreadId = Volatile.Read(ref state.HostThreadId);
@@ -797,14 +580,6 @@ public sealed partial class DirectExecutionBackend
Volatile.Write(ref state.HostThreadId, _prevHostThreadId);
}
TlsSetValue(_backend._hostRspSlotTlsIndex, _prevHostRspSlot);
if (_backend._workerDoneEventTlsIndex != uint.MaxValue)
{
TlsSetValue(_backend._workerDoneEventTlsIndex, 0);
}
if (_backend._tbbAbortEligibleTlsIndex != uint.MaxValue)
{
TlsSetValue(_backend._tbbAbortEligibleTlsIndex, 0);
}
GuestThreadExecution.RestoreGuestThread(_prevGuestThreadHandle);
_activeExecutionBackend = _prevBackend;
_activeCpuContext = _prevContext;
@@ -214,15 +214,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
private nint _guestReturnStub;
private nint _workerAbortStub;
private nint _vehManagedEntryLock;
private uint _workerDoneEventTlsIndex = uint.MaxValue;
private uint _tbbAbortEligibleTlsIndex = uint.MaxValue;
private nint _setEventAddress;
private nint _rawExceptionHandler;
private nint _rawExceptionHandlerStub;
@@ -1050,8 +1041,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
_selfHandlePtr = GCHandle.ToIntPtr(_selfHandle);
_guestTlsBaseTlsIndex = TlsAlloc();
_hostRspSlotTlsIndex = TlsAlloc();
_workerDoneEventTlsIndex = OperatingSystem.IsWindows() ? TlsAlloc() : uint.MaxValue;
_tbbAbortEligibleTlsIndex = OperatingSystem.IsWindows() ? TlsAlloc() : uint.MaxValue;
if (_guestTlsBaseTlsIndex == uint.MaxValue || _hostRspSlotTlsIndex == uint.MaxValue)
{
throw new OutOfMemoryException("Failed to allocate native TLS slots");
@@ -1075,7 +1064,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
{
throw new InvalidOperationException("Failed to resolve kernel32 thread timing functions");
}
_setEventAddress = kernel32 != 0 ? GetProcAddress(kernel32, "SetEvent") : 0;
}
else
{
@@ -1099,29 +1087,13 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
{
throw new OutOfMemoryException("Failed to allocate host stack slot storage");
}
_vehManagedEntryLock = (nint)VirtualAlloc(null, 64u, 12288u, 4u);
if (_vehManagedEntryLock == 0)
{
throw new OutOfMemoryException("Failed to allocate VEH managed-entry lock");
}
// owner (nint) + depth (int); recursive — nested VEH on same thread must reenter.
*(nint*)_vehManagedEntryLock = 0;
*(int*)(_vehManagedEntryLock + sizeof(nint)) = 0;
_unresolvedReturnStub = CreateUnresolvedReturnStub();
_guestReturnStub = CreateGuestReturnStub();
if (_guestReturnStub == 0)
{
throw new OutOfMemoryException("Failed to allocate guest return stub");
}
_workerAbortStub = CreateWorkerAbortStub();
if (_workerAbortStub == 0 && OperatingSystem.IsWindows())
{
Console.Error.WriteLine(
"[LOADER][WARN] Worker abort stub unavailable; TBB execute-fault recover will use host_exit");
}
SetupExceptionHandler();
// Cover the Astro TBB spawn storm (often 812 concurrent tbb_thead).
PrewarmNativeGuestWorkers(Math.Max(NativeWorkerMaxConcurrent, 4));
}
public bool TryExecute(CpuContext context, ulong entryPoint, Generation generation, IReadOnlyDictionary<ulong, string> importStubs, IReadOnlyDictionary<string, ulong> runtimeSymbols, CpuExecutionOptions executionOptions, out OrbisGen2Result result)
@@ -2437,147 +2409,9 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
return (nint)ptr;
}
/// <summary>
/// After a TBB execute-fault, VEH redirects here on a host stack.
/// SetEvent(done) then park forever — ExitThread from VEH CONTINUE_EXECUTION
/// was taking down the whole process (recover logged, no respawning). The
/// renter TerminateThread's the parked worker and respawns a clean loop.
/// </summary>
private unsafe nint CreateWorkerAbortStub()
{
if (!OperatingSystem.IsWindows() ||
_workerDoneEventTlsIndex == uint.MaxValue ||
_tlsGetValueAddress == 0 ||
_setEventAddress == 0)
{
return 0;
}
nint kernel32 = GetModuleHandle("kernel32.dll");
nint getStdHandle = kernel32 != 0 ? GetProcAddress(kernel32, "GetStdHandle") : 0;
nint writeFile = kernel32 != 0 ? GetProcAddress(kernel32, "WriteFile") : 0;
nint flushFileBuffers = kernel32 != 0 ? GetProcAddress(kernel32, "FlushFileBuffers") : 0;
const uint stubSize = 256u;
void* ptr = VirtualAlloc(null, stubSize, 12288u, 4u);
if (ptr == null)
{
return 0;
}
byte* code = (byte*)ptr;
int offset = 0;
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x28); // sub rsp, 0x28
EmitByte(code, ref offset, 0xB9);
EmitUInt32(code, ref offset, _workerDoneEventTlsIndex); // mov ecx, tls
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = _tlsGetValueAddress;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0); // call TlsGetValue
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x85);
EmitByte(code, ref offset, 0xC0); // test rax, rax
EmitByte(code, ref offset, 0x74); EmitByte(code, ref offset, 0x0F); // jz skip SetEvent (15 bytes)
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xC1); // mov rcx, rax
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = _setEventAddress;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0); // call SetEvent
// Breadcrumb on host stack (survives silent teardown better than managed log).
int msgAbsSlot = -1;
if (getStdHandle != 0 && writeFile != 0)
{
ReadOnlySpan<byte> msg = "[LOADER][WARN] tbb_abort_stub SetEvent+park\n"u8;
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x20); // extra shadow for WriteFile args
EmitByte(code, ref offset, 0xB9); EmitUInt32(code, ref offset, unchecked((uint)-12));
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = getStdHandle;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xC1); // mov rcx, handle
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xC3); // mov rbx, handle (nonvolatile for flush)
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
msgAbsSlot = offset;
*(nint*)(code + offset) = 0;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xC2); // mov rdx, msg
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xB8);
EmitUInt32(code, ref offset, (uint)msg.Length);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); // lea r9, [rsp+0x20]
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = writeFile;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
if (flushFileBuffers != 0)
{
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xD9); // mov rcx, rbx
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = flushFileBuffers;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
}
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x20);
}
// Park: do not ExitThread (process-wide silent die after VEH redirect).
int parkOffset = offset;
EmitByte(code, ref offset, 0xF3); EmitByte(code, ref offset, 0x90); // pause
EmitByte(code, ref offset, 0xEB);
EmitByte(code, ref offset, unchecked((byte)(parkOffset - (offset + 1)))); // jmp park
if (msgAbsSlot >= 0)
{
ReadOnlySpan<byte> msgEmbed = "[LOADER][WARN] tbb_abort_stub SetEvent+park\n"u8;
*(nint*)(code + msgAbsSlot) = (nint)ptr + offset;
for (int i = 0; i < msgEmbed.Length; i++)
{
EmitByte(code, ref offset, msgEmbed[i]);
}
}
if (offset > (int)stubSize)
{
Console.Error.WriteLine(
$"[LOADER][ERROR] Worker abort stub overflow: used={offset} cap={stubSize}");
VirtualFree(ptr, 0u, 32768u);
return 0;
}
uint oldProtect = default;
if (!VirtualProtect(ptr, stubSize, 32u, &oldProtect))
{
VirtualFree(ptr, 0u, 32768u);
return 0;
}
FlushInstructionCache(GetCurrentProcess(), ptr, (nuint)offset);
return (nint)ptr;
}
private unsafe nint CreateExceptionHandlerTrampoline(nint managedHandler)
{
// Live VEH trampoline used by SetupExceptionHandler. Must pre-filter
// FastFail / CLR / MSVC C++ / stack-overflow the same way as
// WindowsFaultHandling.CreateHandlerThunk: entering managed VEH while
// the thread is in cooperative GC mode fail-fasts with
// "UnmanagedCallersOnly method from managed code" (tLT1822).
// Extra headroom for native tbb abort + recursive managed-entry spinlock.
const uint stubSize = 2048u;
const uint stubSize = 256u;
void* ptr = VirtualAlloc(null, stubSize, 12288u, 64u);
if (ptr == null)
{
@@ -2586,305 +2420,10 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
byte* code = (byte*)ptr;
int offset = 0;
ReadOnlySpan<uint> nonManagedExceptionCodes =
[
0xE0434352u, // CLR managed exception
0xE06D7363u, // MSVC C++ exception
0xC0000409u, // STATUS_STACK_BUFFER_OVERRUN / FailFast
0xC00000FDu, // STATUS_STACK_OVERFLOW
];
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x01); // mov rax, [rcx]
EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x00); // mov eax, [rax] ExceptionCode
var passJumpOffsets = stackalloc int[nonManagedExceptionCodes.Length];
int fastFailJumpSlot = -1;
for (int i = 0; i < nonManagedExceptionCodes.Length; i++)
{
EmitByte(code, ref offset, 0x3D);
EmitUInt32(code, ref offset, nonManagedExceptionCodes[i]);
EmitByte(code, ref offset, 0x74);
passJumpOffsets[i] = offset;
EmitByte(code, ref offset, 0x00);
if (nonManagedExceptionCodes[i] == 0xC0000409u)
{
fastFailJumpSlot = i;
}
}
EmitByte(code, ref offset, 0xE9); // jmp mainBody (rel32; FastFail breadcrumb sits between)
var mainBodyJumpSlot = offset;
EmitUInt32(code, ref offset, 0u);
int passOffset = offset;
EmitByte(code, ref offset, 0x31); EmitByte(code, ref offset, 0xC0); // xor eax, eax
EmitByte(code, ref offset, 0xC3);
int fastFailPassOffset = offset;
var fastFailLogInstalled = false;
nint kernel32 = GetModuleHandle("kernel32.dll");
nint getStdHandle = kernel32 != 0 ? GetProcAddress(kernel32, "GetStdHandle") : 0;
nint writeFile = kernel32 != 0 ? GetProcAddress(kernel32, "WriteFile") : 0;
if (fastFailJumpSlot >= 0 && getStdHandle != 0 && writeFile != 0)
{
// Prefix + Context.Rip hex (AMD64 CONTEXT.Rip @ 0xF8) + newline.
// Keep in sync with WindowsFaultHandling.CreateHandlerThunk.
ReadOnlySpan<byte> msg =
"[LOADER][FATAL] VEH_PASS FastFail 0xC0000409 (live trampoline; skip managed VEH) rip=0x"u8;
ReadOnlySpan<byte> hexDigits = "0123456789ABCDEF"u8;
// rcx=EXCEPTION_POINTERS*: capture Rip into r10 before clobbering.
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x41);
EmitByte(code, ref offset, 0x08); // mov rax, [rcx+8] ContextRecord*
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x90);
EmitUInt32(code, ref offset, 0xF8u); // mov r10, [rax+0xF8] Rip
EmitByte(code, ref offset, 0x50);
EmitByte(code, ref offset, 0x51);
EmitByte(code, ref offset, 0x52);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x50);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x51);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x52); // push r10 (Rip)
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x40); // sub rsp, 0x40 (hex buf @ +0x30)
EmitByte(code, ref offset, 0xB9); EmitUInt32(code, ref offset, unchecked((uint)-12));
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = getStdHandle;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28); // mov [rsp+0x28], rax stderr handle
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC1);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
var msgAbsSlot = offset;
*(nint*)(code + offset) = 0;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC2);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xB8);
EmitUInt32(code, ref offset, (uint)msg.Length);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x38); EmitUInt32(code, ref offset, 0); // lpOverlapped slot
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = writeFile;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
// Hex-encode Rip. Stack after sub 0x40: [rsp+0x40]=saved Rip (push r10).
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x54); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x40); // mov r10, [rsp+0x40]
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB8);
var hexDigitsAbsSlot = offset;
*(nint*)(code + offset) = 0;
offset += sizeof(nint); // mov r8, hexDigits
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x5C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x30); // lea r11, [rsp+0x30] hex out
EmitByte(code, ref offset, 0xB9); EmitUInt32(code, ref offset, 16u); // ecx = 16 nibbles
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xD0); // mov rax, r10
int hexLoopOffset = offset;
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC1); EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x04); // rol rax, 4
EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC2); // mov edx, eax
EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xE2); EmitByte(code, ref offset, 0x0F); // and edx, 0xF
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0xB6);
EmitByte(code, ref offset, 0x14); EmitByte(code, ref offset, 0x10); // movzx edx, byte [r8+rdx]
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x88); EmitByte(code, ref offset, 0x13); // mov [r11], dl
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xC3); // inc r11
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xC9); // dec ecx
EmitByte(code, ref offset, 0x75);
EmitByte(code, ref offset, unchecked((byte)(hexLoopOffset - (offset + 1)))); // jnz hexLoop (rel8)
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xC6); EmitByte(code, ref offset, 0x03);
EmitByte(code, ref offset, 0x0A); // mov byte [r11], '\n'
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28); // mov rcx, [rsp+0x28] stderr
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x54); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x30); // lea rdx, [rsp+0x30]
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xB8);
EmitUInt32(code, ref offset, 17u); // 16 hex + newline
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x38); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = writeFile;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
// Flush redirected stderr so FastFail rip survives process teardown.
nint flushFileBuffers = kernel32 != 0 ? GetProcAddress(kernel32, "FlushFileBuffers") : 0;
if (flushFileBuffers != 0)
{
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28); // mov rcx, stderr
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = flushFileBuffers;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
}
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x40);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x5A); // pop r10
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x59);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x58);
EmitByte(code, ref offset, 0x5A);
EmitByte(code, ref offset, 0x59);
EmitByte(code, ref offset, 0x58);
EmitByte(code, ref offset, 0x31); EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0xC3);
var msgOffset = offset;
for (int i = 0; i < msg.Length; i++)
{
EmitByte(code, ref offset, msg[i]);
}
var hexDigitsOffset = offset;
for (int i = 0; i < hexDigits.Length; i++)
{
EmitByte(code, ref offset, hexDigits[i]);
}
*(nint*)(code + msgAbsSlot) = (nint)ptr + msgOffset;
*(nint*)(code + hexDigitsAbsSlot) = (nint)ptr + hexDigitsOffset;
code[passJumpOffsets[fastFailJumpSlot]] =
checked((byte)(fastFailPassOffset - (passJumpOffsets[fastFailJumpSlot] + 1)));
fastFailLogInstalled = true;
}
int mainBodyOffset = offset;
*(int*)(code + mainBodyJumpSlot) = mainBodyOffset - (mainBodyJumpSlot + sizeof(int));
for (int i = 0; i < nonManagedExceptionCodes.Length; i++)
{
if (i == fastFailJumpSlot && fastFailLogInstalled)
{
continue;
}
code[passJumpOffsets[i]] = checked((byte)(passOffset - (passJumpOffsets[i] + 1)));
}
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x54); // push r12
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x55); // push r13
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xE4); // mov r12, rsp
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xCD); // mov r13, rcx
// Native worker EXECUTE-AV abort without managed VEH.
// Do NOT catch read/write AVs — workers need managed lazy-commit (tLTJ
// silent-die when every worker AV was aborted). Execute faults on
// tbb_thead are the concurrent-managed FailFast case (tLTC).
int tbbFallthroughJump = -1;
if (_workerAbortStub != 0 &&
_tlsGetValueAddress != 0 &&
_hostRspSlotTlsIndex != uint.MaxValue)
{
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x45); EmitByte(code, ref offset, 0x00); // mov rax, [r13]
EmitByte(code, ref offset, 0x81); EmitByte(code, ref offset, 0x38);
EmitUInt32(code, ref offset, 0xC0000005u); // cmp dword [rax], AV
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
tbbFallthroughJump = offset;
EmitUInt32(code, ref offset, 0u); // jne fallthrough
// ExceptionInformation[0] == 8 → EXECUTE (DEP). Offset 32 on x64 RECORD.
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xB8); EmitUInt32(code, ref offset, 32u);
EmitByte(code, ref offset, 0x08); // cmp qword [rax+32], 8
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
var tbbNotExecuteJump = offset;
EmitUInt32(code, ref offset, 0u);
int tbbNotEligibleJump = -1;
if (_tbbAbortEligibleTlsIndex != uint.MaxValue)
{
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0xB9);
EmitUInt32(code, ref offset, _tbbAbortEligibleTlsIndex);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = _tlsGetValueAddress;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x85);
EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x84);
tbbNotEligibleJump = offset;
EmitUInt32(code, ref offset, 0u);
}
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0xB9);
EmitUInt32(code, ref offset, _hostRspSlotTlsIndex);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = _tlsGetValueAddress;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x85);
EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x84);
var tbbNoHostRspJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x00); // mov r8, [rax] hostRsp
EmitByte(code, ref offset, 0x4D); EmitByte(code, ref offset, 0x85);
EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x84);
var tbbZeroRspJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xE0); EmitByte(code, ref offset, 0xF0); // and r8, ~0xF
EmitByte(code, ref offset, 0x4D); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x4D); EmitByte(code, ref offset, 0x08); // mov r9, [r13+8]
EmitByte(code, ref offset, 0x4D); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0x81); EmitUInt32(code, ref offset, 0x98u); // Context.Rsp
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = _workerAbortStub;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0x81); EmitUInt32(code, ref offset, 0xF8u); // Context.Rip
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x78);
EmitUInt32(code, ref offset, 0u); // Context.Rax = 0
EmitByte(code, ref offset, 0xB8); EmitUInt32(code, ref offset, unchecked((uint)-1));
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0xE4); // mov rsp, r12
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x5D);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x5C);
EmitByte(code, ref offset, 0xC3);
int tbbFallthroughOffset = offset;
*(int*)(code + tbbFallthroughJump) = tbbFallthroughOffset - (tbbFallthroughJump + sizeof(int));
*(int*)(code + tbbNotExecuteJump) = tbbFallthroughOffset - (tbbNotExecuteJump + sizeof(int));
if (tbbNotEligibleJump >= 0)
{
*(int*)(code + tbbNotEligibleJump) = tbbFallthroughOffset - (tbbNotEligibleJump + sizeof(int));
}
*(int*)(code + tbbNoHostRspJump) = tbbFallthroughOffset - (tbbNoHostRspJump + sizeof(int));
*(int*)(code + tbbZeroRspJump) = tbbFallthroughOffset - (tbbZeroRspJump + sizeof(int));
}
EmitByte(code, ref offset, 0x65); EmitByte(code, ref offset, 0x48); // mov rax, gs:[8]
EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x04); EmitByte(code, ref offset, 0x25);
EmitUInt32(code, ref offset, 8u);
@@ -2901,67 +2440,11 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x28);
// Serialize managed VEH entry (recursive spinlock). Concurrent UnmanagedCallersOnly
// FailFast was the tLTQ silent mid-TBB pattern (enter without abort breadcrumb).
// Lock layout: [0]=owner UniqueThread (nint), [8]=depth (int).
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB9);
*(nint*)(code + offset) = _vehManagedEntryLock;
offset += sizeof(nint); // mov r9, lock*
EmitByte(code, ref offset, 0x65); EmitByte(code, ref offset, 0x4C);
EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x14);
EmitByte(code, ref offset, 0x25); EmitUInt32(code, ref offset, 0x48u); // mov r10, gs:[0x48]
int hostAcquireSpin = offset;
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x01); // mov rax, [r9]
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x39); EmitByte(code, ref offset, 0xD0); // cmp rax, r10
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x84);
int hostMineJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x85); EmitByte(code, ref offset, 0xC0); // test rax, rax
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int hostPauseJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0xF0); EmitByte(code, ref offset, 0x4C);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0xB1); EmitByte(code, ref offset, 0x11); // lock cmpxchg [r9], r10
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int hostRetryJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x08);
EmitUInt32(code, ref offset, 1u); // mov dword [r9+8], 1
EmitByte(code, ref offset, 0xE9);
int hostGotJump = offset;
EmitUInt32(code, ref offset, 0u);
int hostPauseOffset = offset;
EmitByte(code, ref offset, 0xF3); EmitByte(code, ref offset, 0x90); // pause
EmitByte(code, ref offset, 0xE9);
int hostPauseBackJump = offset;
EmitUInt32(code, ref offset, 0u);
int hostMineOffset = offset;
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xFF);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x08); // inc dword [r9+8]
int hostGotOffset = offset;
*(int*)(code + hostMineJump) = hostMineOffset - (hostMineJump + sizeof(int));
*(int*)(code + hostPauseJump) = hostPauseOffset - (hostPauseJump + sizeof(int));
*(int*)(code + hostRetryJump) = hostAcquireSpin - (hostRetryJump + sizeof(int));
*(int*)(code + hostGotJump) = hostGotOffset - (hostGotJump + sizeof(int));
*(int*)(code + hostPauseBackJump) = hostAcquireSpin - (hostPauseBackJump + sizeof(int));
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xE9); // mov rcx, r13
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = managedHandler;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB9);
*(nint*)(code + offset) = _vehManagedEntryLock;
offset += sizeof(nint); // mov r9, lock*
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xFF);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x08); // dec dword [r9+8]
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int hostStillJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x01); EmitUInt32(code, ref offset, 0u); // mov qword [r9], 0
int hostStillOffset = offset;
*(int*)(code + hostStillJump) = hostStillOffset - (hostStillJump + sizeof(int));
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0xE9);
int hostRestoreJump = offset;
@@ -2987,64 +2470,11 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xDC); // mov rsp, r11
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB9);
*(nint*)(code + offset) = _vehManagedEntryLock;
offset += sizeof(nint); // mov r9, lock*
EmitByte(code, ref offset, 0x65); EmitByte(code, ref offset, 0x4C);
EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x14);
EmitByte(code, ref offset, 0x25); EmitUInt32(code, ref offset, 0x48u); // mov r10, gs:[0x48]
int guestAcquireSpin = offset;
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x01); // mov rax, [r9]
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x39); EmitByte(code, ref offset, 0xD0); // cmp rax, r10
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x84);
int guestMineJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x85); EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int guestPauseJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0xF0); EmitByte(code, ref offset, 0x4C);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0xB1); EmitByte(code, ref offset, 0x11);
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int guestRetryJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x08);
EmitUInt32(code, ref offset, 1u);
EmitByte(code, ref offset, 0xE9);
int guestGotJump = offset;
EmitUInt32(code, ref offset, 0u);
int guestPauseOffset = offset;
EmitByte(code, ref offset, 0xF3); EmitByte(code, ref offset, 0x90);
EmitByte(code, ref offset, 0xE9);
int guestPauseBackJump = offset;
EmitUInt32(code, ref offset, 0u);
int guestMineOffset = offset;
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xFF);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x08);
int guestGotOffset = offset;
*(int*)(code + guestMineJump) = guestMineOffset - (guestMineJump + sizeof(int));
*(int*)(code + guestPauseJump) = guestPauseOffset - (guestPauseJump + sizeof(int));
*(int*)(code + guestRetryJump) = guestAcquireSpin - (guestRetryJump + sizeof(int));
*(int*)(code + guestGotJump) = guestGotOffset - (guestGotJump + sizeof(int));
*(int*)(code + guestPauseBackJump) = guestAcquireSpin - (guestPauseBackJump + sizeof(int));
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xE9); // mov rcx, r13
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = managedHandler;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB9);
*(nint*)(code + offset) = _vehManagedEntryLock;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xFF);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0x08); // dec dword [r9+8]
EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0x85);
int guestStillJump = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x01); EmitUInt32(code, ref offset, 0u);
int guestStillOffset = offset;
*(int*)(code + guestStillJump) = guestStillOffset - (guestStillJump + sizeof(int));
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0xE9);
int guestRestoreJump = offset;
@@ -3065,18 +2495,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
*(int*)(code + missingHostStackJump) = passThroughOffset - (missingHostStackJump + sizeof(int));
*(int*)(code + guestRestoreJump) = restoreOffset - (guestRestoreJump + sizeof(int));
if (offset > (int)stubSize)
{
Console.Error.WriteLine(
$"[LOADER][ERROR] Exception handler trampoline overflow: used={offset} cap={stubSize}");
VirtualFree(ptr, 0, 0x8000u);
return 0;
}
Console.Error.WriteLine(
$"[LOADER][INFO] VEH trampoline built: bytes={offset} native_worker_abort=" +
$"{(_workerAbortStub != 0 && _hostRspSlotTlsIndex != uint.MaxValue)}");
uint oldProtect = default;
VirtualProtect(ptr, stubSize, 32u, &oldProtect);
FlushInstructionCache(GetCurrentProcess(), ptr, (nuint)offset);
@@ -5731,27 +5149,16 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
return GuestNativeCallExitReason.Exception;
}
FlushInstructionCache(GetCurrentProcess(), ptr, stubSize);
if (!TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot))
{
reason = "failed to bind host-RSP storage for guest thread stub";
return GuestNativeCallExitReason.Exception;
}
ActiveGuestThreadYieldRequested = false;
ActiveGuestThreadYieldReason = null;
try
{
// TBB execute-AV recover needs native-worker TLS (eligible/done).
// Other guests stay on CallNativeEntry — full native-worker migration
// increased splash hangs / UnmanagedCallersOnly (tLTN/tLTO).
int nativeReturn;
if (name == "tbb_thead")
{
nativeReturn = RunGuestEntryStub(ptr, hostRspSlot, requireNativeWorker: true);
}
else
{
if (!TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot))
{
reason = "failed to bind host-RSP storage for guest thread stub";
return GuestNativeCallExitReason.Exception;
}
nativeReturn = CallNativeEntry(ptr);
}
var nativeReturn = CallNativeEntry(ptr);
if (ActiveGuestThreadYieldRequested)
{
reason = ActiveGuestThreadYieldReason ?? "guest thread blocked";
@@ -5897,24 +5304,16 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
return GuestNativeCallExitReason.Exception;
}
FlushInstructionCache(GetCurrentProcess(), ptr, stubSize);
if (!TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot))
{
reason = "failed to bind host-RSP storage for guest continuation stub";
return GuestNativeCallExitReason.Exception;
}
ActiveGuestThreadYieldRequested = false;
ActiveGuestThreadYieldReason = null;
try
{
int nativeReturn;
if (name == "tbb_thead")
{
nativeReturn = RunGuestEntryStub(ptr, hostRspSlot, requireNativeWorker: true);
}
else
{
if (!TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot))
{
reason = "failed to bind host-RSP storage for guest continuation stub";
return GuestNativeCallExitReason.Exception;
}
nativeReturn = CallNativeEntry(ptr);
}
var nativeReturn = CallNativeEntry(ptr);
if (ActiveGuestThreadYieldRequested)
{
reason = ActiveGuestThreadYieldReason ?? "guest thread blocked";
@@ -7047,16 +6446,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
VirtualFree((void*)_hostRspSlotStorage, 0u, 32768u);
_hostRspSlotStorage = 0;
}
if (_vehManagedEntryLock != 0)
{
VirtualFree((void*)_vehManagedEntryLock, 0u, 32768u);
_vehManagedEntryLock = 0;
}
if (_workerAbortStack != 0)
{
VirtualFree((void*)_workerAbortStack, 0u, 32768u);
_workerAbortStack = 0;
}
if (_guestTlsBaseTlsIndex != uint.MaxValue)
{
TlsFree(_guestTlsBaseTlsIndex);
@@ -7067,11 +6456,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
TlsFree(_hostRspSlotTlsIndex);
_hostRspSlotTlsIndex = uint.MaxValue;
}
if (_workerDoneEventTlsIndex != uint.MaxValue)
{
TlsFree(_workerDoneEventTlsIndex);
_workerDoneEventTlsIndex = uint.MaxValue;
}
if (_unresolvedReturnStub != 0)
{
VirtualFree((void*)_unresolvedReturnStub, 0u, 32768u);
@@ -7082,11 +6466,6 @@ public sealed unsafe partial class DirectExecutionBackend : INativeCpuBackend, I
VirtualFree((void*)_guestReturnStub, 0u, 32768u);
_guestReturnStub = 0;
}
if (_workerAbortStub != 0)
{
VirtualFree((void*)_workerAbortStub, 0u, 32768u);
_workerAbortStub = 0;
}
if (_guestContextTransferStub != 0)
{
VirtualFree((void*)_guestContextTransferStub, 0u, 32768u);
@@ -24,7 +24,7 @@ internal sealed unsafe partial class WindowsFaultHandling : IHostFaultHandling
public nint CreateHandlerThunk(nint managedCallback, uint hostRspSwitchTlsSlot, nint tlsGetValueAddress)
{
const uint stubSize = 1024u;
const uint stubSize = 256u;
void* ptr = (void*)_memory.Allocate(0, stubSize, HostPageProtection.ReadWriteExecute);
if (ptr == null)
{
@@ -43,15 +43,11 @@ internal sealed unsafe partial class WindowsFaultHandling : IHostFaultHandling
// managed code; the CLR's own VEH handles its exceptions. MSVC C++ exceptions
// (Vulkan drivers, host CRT) are excluded too: the managed handler only ever
// returned CONTINUE_SEARCH for them.
//
// FastFail (0xC0000409) is logged from this native path only: managed VEH never
// sees it (tLT1821 silent exits after TBB AV recovery).
ReadOnlySpan<uint> nonManagedExceptionCodes =
[WindowsFaultCodes.ClrManagedException, 0xE06D7363u, WindowsFaultCodes.FastFail, WindowsFaultCodes.StackOverflow];
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x01); // mov rax, [rcx] (ExceptionRecord*)
EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x00); // mov eax, [rax] (ExceptionCode)
var passJumpOffsets = stackalloc int[nonManagedExceptionCodes.Length];
int fastFailJumpSlot = -1;
for (int i = 0; i < nonManagedExceptionCodes.Length; i++)
{
EmitByte(code, ref offset, 0x3D); // cmp eax, imm32
@@ -59,162 +55,13 @@ internal sealed unsafe partial class WindowsFaultHandling : IHostFaultHandling
EmitByte(code, ref offset, 0x74); // je pass
passJumpOffsets[i] = offset;
EmitByte(code, ref offset, 0x00);
if (nonManagedExceptionCodes[i] == WindowsFaultCodes.FastFail)
{
fastFailJumpSlot = i;
}
}
EmitByte(code, ref offset, 0xE9); // jmp mainBody rel32 (FastFail breadcrumb sits between)
var mainBodyJumpSlot = offset;
EmitUInt32(code, ref offset, 0u);
EmitByte(code, ref offset, 0xEB); EmitByte(code, ref offset, 0x03); // jmp over pass block
int passOffset = offset;
EmitByte(code, ref offset, 0x31); EmitByte(code, ref offset, 0xC0); // pass: xor eax, eax (EXCEPTION_CONTINUE_SEARCH)
EmitByte(code, ref offset, 0xC3); // ret
// FastFail: native stderr breadcrumb with Context.Rip (no managed entry), then CONTINUE_SEARCH.
// Keep in sync with DirectExecutionBackend.CreateExceptionHandlerTrampoline.
int fastFailPassOffset = offset;
var fastFailLogInstalled = false;
if (fastFailJumpSlot >= 0 &&
NativeLibrary.TryLoad("kernel32.dll", out var kernel32) &&
NativeLibrary.TryGetExport(kernel32, "GetStdHandle", out var getStdHandle) &&
NativeLibrary.TryGetExport(kernel32, "WriteFile", out var writeFile))
{
ReadOnlySpan<byte> msg =
"[LOADER][FATAL] VEH_PASS FastFail 0xC0000409 (native; no managed VEH) rip=0x"u8;
ReadOnlySpan<byte> hexDigits = "0123456789ABCDEF"u8;
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x41);
EmitByte(code, ref offset, 0x08); // mov rax, [rcx+8]
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8B); EmitByte(code, ref offset, 0x90);
EmitUInt32(code, ref offset, 0xF8u); // mov r10, [rax+0xF8]
EmitByte(code, ref offset, 0x50); // push rax
EmitByte(code, ref offset, 0x51); // push rcx
EmitByte(code, ref offset, 0x52); // push rdx
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x50); // push r8
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x51); // push r9
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x52); // push r10
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xEC); EmitByte(code, ref offset, 0x40); // sub rsp, 0x40
EmitByte(code, ref offset, 0xB9); EmitUInt32(code, ref offset, unchecked((uint)-12));
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = getStdHandle;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28); // mov [rsp+0x28], rax
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC1);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
var msgAbsSlot = offset;
*(nint*)(code + offset) = 0;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC2);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xB8);
EmitUInt32(code, ref offset, (uint)msg.Length);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x38); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = writeFile;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x54); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x40); // mov r10, [rsp+0x40]
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xB8);
var hexDigitsAbsSlot = offset;
*(nint*)(code + offset) = 0;
offset += sizeof(nint);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x5C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x30); // lea r11, [rsp+0x30]
EmitByte(code, ref offset, 0xB9); EmitUInt32(code, ref offset, 16u);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xD0);
int hexLoopOffset = offset;
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC1); EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0x04);
EmitByte(code, ref offset, 0x89); EmitByte(code, ref offset, 0xC2);
EmitByte(code, ref offset, 0x83); EmitByte(code, ref offset, 0xE2); EmitByte(code, ref offset, 0x0F);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x0F); EmitByte(code, ref offset, 0xB6);
EmitByte(code, ref offset, 0x14); EmitByte(code, ref offset, 0x10);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x88); EmitByte(code, ref offset, 0x13);
EmitByte(code, ref offset, 0x49); EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xC3);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xC9);
EmitByte(code, ref offset, 0x75);
EmitByte(code, ref offset, unchecked((byte)(hexLoopOffset - (offset + 1)))); // jnz rel8
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xC6); EmitByte(code, ref offset, 0x03);
EmitByte(code, ref offset, 0x0A);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8B);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x28);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x54); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x30);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0xB8);
EmitUInt32(code, ref offset, 17u);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x8D);
EmitByte(code, ref offset, 0x4C); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x20); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xC7);
EmitByte(code, ref offset, 0x44); EmitByte(code, ref offset, 0x24);
EmitByte(code, ref offset, 0x38); EmitUInt32(code, ref offset, 0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0xB8);
*(nint*)(code + offset) = writeFile;
offset += sizeof(nint);
EmitByte(code, ref offset, 0xFF); EmitByte(code, ref offset, 0xD0);
EmitByte(code, ref offset, 0x48); EmitByte(code, ref offset, 0x83);
EmitByte(code, ref offset, 0xC4); EmitByte(code, ref offset, 0x40);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x5A);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x59);
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x58);
EmitByte(code, ref offset, 0x5A);
EmitByte(code, ref offset, 0x59);
EmitByte(code, ref offset, 0x58);
EmitByte(code, ref offset, 0x31); EmitByte(code, ref offset, 0xC0);
EmitByte(code, ref offset, 0xC3);
var msgOffset = offset;
for (int i = 0; i < msg.Length; i++)
{
EmitByte(code, ref offset, msg[i]);
}
var hexDigitsOffset = offset;
for (int i = 0; i < hexDigits.Length; i++)
{
EmitByte(code, ref offset, hexDigits[i]);
}
*(nint*)(code + msgAbsSlot) = (nint)ptr + msgOffset;
*(nint*)(code + hexDigitsAbsSlot) = (nint)ptr + hexDigitsOffset;
code[passJumpOffsets[fastFailJumpSlot]] =
checked((byte)(fastFailPassOffset - (passJumpOffsets[fastFailJumpSlot] + 1)));
fastFailLogInstalled = true;
}
int mainBodyOffset = offset;
*(int*)(code + mainBodyJumpSlot) = mainBodyOffset - (mainBodyJumpSlot + sizeof(int));
for (int i = 0; i < nonManagedExceptionCodes.Length; i++)
{
if (i == fastFailJumpSlot && fastFailLogInstalled)
{
continue;
}
code[passJumpOffsets[i]] = checked((byte)(passOffset - (passJumpOffsets[i] + 1)));
}
EmitByte(code, ref offset, 0x41); EmitByte(code, ref offset, 0x54); // push r12
@@ -238,15 +238,7 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
var alignedSize = (size + 0xFFF) & ~0xFFFUL;
var protection = executable ? PAGE_EXECUTE_READWRITE : PAGE_READWRITE;
var hostProtection = executable ? HostPageProtection.ReadWriteExecute : HostPageProtection.ReadWrite;
// Reserve address space only for very large non-executable regions; commit is done lazily later.
var reservedOnly = !executable &&
alignedSize >= LargeDataReserveThreshold &&
alignedSize > FullCommitRegionLimit;
var result = reservedOnly
? _hostMemory.Reserve(desiredAddress, alignedSize, HostPageProtection.ReadWrite)
: _hostMemory.Allocate(desiredAddress, alignedSize, hostProtection);
var result = _hostMemory.Allocate(desiredAddress, alignedSize, hostProtection);
if (result == 0)
{
return false;
@@ -260,8 +252,6 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
return false;
}
var state = reservedOnly ? ReserveRegion(actualAddress, alignedSize) : "n/a";
_gate.EnterWriteLock();
try
{
@@ -270,7 +260,7 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
VirtualAddress = actualAddress,
Size = alignedSize,
IsExecutable = executable,
IsReservedOnly = reservedOnly,
IsReservedOnly = false,
Protection = protection
});
}
@@ -279,7 +269,6 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
_gate.ExitWriteLock();
}
var allocationKind = executable ? "executable memory" : "data memory";
TraceVmem($"Allocated exact {allocationKind}: 0x{actualAddress:X16} - 0x{actualAddress + alignedSize:X16} ({alignedSize} bytes)");
return true;
@@ -372,7 +361,44 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
var actualAddress = result;
var lazyPrimeState = reservedOnly ? ReserveRegion(actualAddress, alignedSize) : "n/a";
var lazyPrimeState = "n/a";
if (reservedOnly)
{
var primeBytes = Math.Min(alignedSize, LazyReservePrimeBytes);
if (primeBytes != 0)
{
ulong committedBytes = 0;
while (committedBytes < primeBytes)
{
var remaining = primeBytes - committedBytes;
var chunkBytes = Math.Min(remaining, LazyReservePrimeChunkBytes);
var commitAddress = actualAddress + committedBytes;
if (!_hostMemory.Commit(commitAddress, chunkBytes, HostPageProtection.ReadWrite))
{
break;
}
committedBytes += chunkBytes;
}
if (committedBytes != 0)
{
lazyPrimeState = committedBytes == primeBytes
? $"ok:{committedBytes:X}"
: $"partial:{committedBytes:X}/{primeBytes:X}";
TraceVmem($"Primed lazy region: 0x{actualAddress:X16} - 0x{actualAddress + committedBytes:X16} ({committedBytes} bytes)");
}
else
{
lazyPrimeState = $"fail:{primeBytes:X}";
TraceVmem($"Failed to prime lazy region at 0x{actualAddress:X16} ({primeBytes} bytes), continuing with on-demand commit");
}
}
else
{
lazyPrimeState = "skip:0";
}
}
_gate.EnterWriteLock();
try
@@ -399,41 +425,6 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
return actualAddress;
}
private string ReserveRegion(ulong actualAddress, ulong alignedSize)
{
var primeBytes = Math.Min(alignedSize, LazyReservePrimeBytes);
if (primeBytes == 0)
{
return "skip:0";
}
ulong committedBytes = 0;
while (committedBytes < primeBytes)
{
var remaining = primeBytes - committedBytes;
var chunkBytes = Math.Min(remaining, LazyReservePrimeChunkBytes);
var commitAddress = actualAddress + committedBytes;
if (!_hostMemory.Commit(commitAddress, chunkBytes, HostPageProtection.ReadWrite))
{
break;
}
committedBytes += chunkBytes;
}
if (committedBytes != 0)
{
var state = committedBytes == primeBytes
? $"ok:{committedBytes:X}"
: $"partial:{committedBytes:X}/{primeBytes:X}";
TraceVmem($"region: 0x{actualAddress:X16} - 0x{actualAddress + committedBytes:X16} ({committedBytes} bytes)");
return state;
}
TraceVmem($"Failed to reserve region at 0x{actualAddress:X16} ({primeBytes} bytes)!");
return $"fail:{primeBytes:X}";
}
public bool TryBackFixedRange(ulong address, ulong size, bool executable)
{
if (size == 0)
-23
View File
@@ -221,29 +221,6 @@ public static class GuestThreadExecution
public static IGuestThreadScheduler? Scheduler { get; set; }
/// <summary>
/// Fired when a guest thread is torn down without a clean pthread_exit
/// (e.g. TBB execute-AV → worker_abort). Libs use this to abandon mutexes.
/// </summary>
public static event Func<ulong, string, int>? GuestThreadAbandoned;
public static int NotifyGuestThreadAbandoned(ulong threadHandle, string reason)
{
if (threadHandle == 0 || GuestThreadAbandoned is null)
{
return 0;
}
try
{
return GuestThreadAbandoned.Invoke(threadHandle, reason);
}
catch
{
return 0;
}
}
public static bool IsGuestThread => _currentGuestThreadHandle != 0;
public static ulong CurrentGuestThreadHandle => _currentGuestThreadHandle;
File diff suppressed because it is too large Load Diff
-184
View File
@@ -6,55 +6,6 @@ using System.Runtime.CompilerServices;
namespace SharpEmu.Libs.Agc;
/// <summary>Which in-block address equation a <see cref="DetileParams"/> carries.</summary>
internal enum DetileEquation
{
/// <summary>Unsupported mode/format; caller must use the CPU path or raw upload.</summary>
None,
/// <summary>Exact AddrLib XOR equation (RDNA2 modes 5/9/24/27): factored X/Y terms.</summary>
ExactXor,
/// <summary>Other modes: a precomputed in-block Morton/standard element-offset table.</summary>
BlockTable,
}
/// <summary>
/// Backend-agnostic description of how to deswizzle one surface, produced by
/// <see cref="GnmTiling.GetDetileParams"/>. 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 <see cref="GnmTiling.TryDetile"/> fallback. The single
/// shared addressing formula both consume is:
/// <code>
/// inBlockByte = Equation == ExactXor
/// ? XByteTerm[x &amp; XMask] ^ YByteTerm[y &amp; YMask]
/// : BlockTable[(y % BlockHeight) * BlockWidth + (x % BlockWidth)] * BytesPerElement;
/// srcByte = ((y / BlockHeight) * BlocksPerRow + (x / BlockWidth)) * BlockBytes + inBlockByte;
/// </code>
/// </summary>
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)
{
/// <summary>False when the mode/format is not GPU-portable (Equation == None).</summary>
public bool IsSupported => Equation != DetileEquation.None;
}
/// <summary>
/// Deswizzles RDNA2 (GFX10) tiled texture surfaces into linear layout so they
/// can be uploaded to Vulkan. PS5 stores most textures in a swizzled layout
@@ -550,141 +501,6 @@ internal static unsafe class GnmTiling
return true;
}
/// <summary>
/// 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
/// <see cref="DetileParams.IsSupported"/> == false (Equation == None) when the
/// mode/format is not GPU-portable, so the caller keeps the CPU
/// <see cref="TryDetile"/> path or a raw upload. Reuses the same helpers and
/// caches as <see cref="TryDetile"/>, so the two never disagree on addressing.
/// </summary>
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);
}
/// <summary>
/// CPU deswizzle driven entirely by a resolved <see cref="DetileParams"/> — the
/// exact addressing the Vulkan/Metal compute kernel runs per texel, so a
/// backend that packaged <paramref name="parameters"/> for the GPU path can
/// fall back to this without re-deriving the swizzle. Copies
/// <c>ElementsWide * ElementsHigh</c> elements from <paramref name="tiled"/>
/// into <paramref name="linear"/>; 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 <paramref name="tiled"/> degrades gracefully.
/// </summary>
public static bool DetileWithParams(
in DetileParams parameters,
ReadOnlySpan<byte> tiled,
Span<byte> 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,
-241
View File
@@ -21,7 +21,6 @@ public static class AjmExports
private const int MaxInstanceIndex = 0x2FFF;
private static readonly ConcurrentDictionary<uint, AjmContextState> Contexts = new();
private static int _nextContextId;
private static int _nextBatchId;
private sealed class AjmContextState
{
@@ -228,250 +227,10 @@ public static class AjmExports
return 0;
}
/// <summary>
/// Enqueues a decode job on a batch. Titles call this on the Bink/AJM hot
/// path; leaving it unresolved floods Import WARN spam. This is a silence
/// stub, not a codec: advance the batch cursor and report the input as
/// consumed with silence produced so the title does not spin on the same
/// packet.
/// </summary>
[SysAbiExport(
Nid = "39WxhR-ePew",
ExportName = "sceAjmBatchJobDecode",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceAjm")]
public static int AjmBatchJobDecode(CpuContext ctx)
{
var infoAddress = ctx[CpuRegister.Rdi];
var instanceId = unchecked((uint)ctx[CpuRegister.Rsi]);
var inputAddress = ctx[CpuRegister.Rdx];
var inputSize = ctx[CpuRegister.Rcx];
var outputAddress = ctx[CpuRegister.R8];
var outputSize = ctx[CpuRegister.R9];
var resultAddress = ReadStackArg64(ctx, 0);
if (infoAddress == 0)
{
return ctx.SetReturn(OrbisAjmErrorInvalidParameter);
}
// Best-effort: bump the batch cursor when the guest filled AjmBatchInfo.
// Still succeed without it — the unresolved stub returned 0 and titles
// keep calling; failing here would reintroduce hot-path spam via retries.
_ = TryAppendBatchJob(ctx, infoAddress, AjmJobRunSize);
// Silence: clear PCM out and claim full input consumed so the guest
// advances its bitstream cursor instead of re-submitting forever.
if (outputAddress != 0 && outputSize != 0 && outputSize <= MaxSilentPcmBytes)
{
ClearGuestMemory(ctx, outputAddress, outputSize);
}
WriteDecodeStreamResult(
ctx,
resultAddress,
inputConsumed: inputSize > int.MaxValue ? int.MaxValue : (int)inputSize,
outputWritten: 0,
totalDecodedSamples: 0,
frames: inputSize != 0 || outputSize != 0 ? 1u : 0u);
Trace(
$"batch_job_decode info=0x{infoAddress:X16} instance=0x{instanceId:X8} " +
$"in=0x{inputAddress:X16}+0x{inputSize:X} out=0x{outputAddress:X16}+0x{outputSize:X} " +
$"result=0x{resultAddress:X16}");
return ctx.SetReturn(0);
}
/// <summary>
/// Submits a built batch. Instant-complete silence stub: publish a batch id
/// and clear any error out. Decode sidebands were already filled at
/// job-enqueue time.
/// </summary>
[SysAbiExport(
Nid = "5tOfnaClcqM",
ExportName = "sceAjmBatchStart",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceAjm")]
public static int AjmBatchStart(CpuContext ctx)
{
var contextId = unchecked((uint)ctx[CpuRegister.Rdi]);
var infoAddress = ctx[CpuRegister.Rsi];
var priority = unchecked((int)ctx[CpuRegister.Rdx]);
var errorAddress = ctx[CpuRegister.Rcx];
var batchOutAddress = ctx[CpuRegister.R8];
if (infoAddress == 0 || batchOutAddress == 0)
{
return ctx.SetReturn(OrbisAjmErrorInvalidParameter);
}
ClearAjmBatchError(ctx, errorAddress);
var batchId = unchecked((uint)Interlocked.Increment(ref _nextBatchId));
Span<byte> batchValue = stackalloc byte[sizeof(uint)];
BinaryPrimitives.WriteUInt32LittleEndian(batchValue, batchId);
if (!ctx.Memory.TryWrite(batchOutAddress, batchValue))
{
return ctx.SetReturn(OrbisAjmErrorInvalidParameter);
}
Trace(
$"batch_start context={contextId} info=0x{infoAddress:X16} " +
$"priority={priority} batch={batchId} error=0x{errorAddress:X16}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "-qLsfDAywIY",
ExportName = "sceAjmBatchWait",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceAjm")]
public static int AjmBatchWait(CpuContext ctx)
{
// Batches complete synchronously in Start; Wait is a no-op success.
var errorAddress = ctx[CpuRegister.Rcx];
ClearAjmBatchError(ctx, errorAddress);
Trace(
$"batch_wait context={unchecked((uint)ctx[CpuRegister.Rdi])} " +
$"batch={unchecked((uint)ctx[CpuRegister.Rsi])} " +
$"timeout={unchecked((uint)ctx[CpuRegister.Rdx])}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "NVDXiUesSbA",
ExportName = "sceAjmBatchCancel",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceAjm")]
public static int AjmBatchCancel(CpuContext ctx)
{
Trace(
$"batch_cancel context={unchecked((uint)ctx[CpuRegister.Rdi])} " +
$"batch={unchecked((uint)ctx[CpuRegister.Rsi])}");
return ctx.SetReturn(0);
}
internal static void ResetForTests()
{
Contexts.Clear();
Interlocked.Exchange(ref _nextContextId, 0);
Interlocked.Exchange(ref _nextBatchId, 0);
}
// AjmBatchInfo: buffer, offset, size, last_good_job, last_good_job_ra (5× u64).
private const ulong AjmBatchInfoOffsetField = 8;
private const ulong AjmBatchInfoSizeField = 16;
private const ulong AjmBatchInfoLastGoodJobField = 24;
private const ulong AjmJobRunSize = 64;
private const ulong MaxSilentPcmBytes = 1 << 20;
// AjmSidebandResult (8) + AjmSidebandStream (16) + AjmSidebandMFrame (8).
private const int DecodeSidebandBytes = 32;
private static bool TryAppendBatchJob(CpuContext ctx, ulong infoAddress, ulong jobSize)
{
if (!TryReadUInt64(ctx, infoAddress, out var buffer) ||
!TryReadUInt64(ctx, infoAddress + AjmBatchInfoOffsetField, out var offset) ||
!TryReadUInt64(ctx, infoAddress + AjmBatchInfoSizeField, out var size))
{
return false;
}
if (buffer == 0 || jobSize == 0 || offset > size || size - offset < jobSize)
{
return false;
}
var jobAddress = buffer + offset;
ClearGuestMemory(ctx, jobAddress, jobSize);
return TryWriteUInt64(ctx, infoAddress + AjmBatchInfoLastGoodJobField, jobAddress) &&
TryWriteUInt64(ctx, infoAddress + AjmBatchInfoOffsetField, offset + jobSize);
}
// AjmBatchError: int error_code; const void* job_addr; uint32_t cmd_offset; const void* job_ra;
private const int AjmBatchErrorBytes = 24;
private static void ClearAjmBatchError(CpuContext ctx, ulong errorAddress)
{
if (errorAddress == 0)
{
return;
}
Span<byte> error = stackalloc byte[AjmBatchErrorBytes];
error.Clear();
_ = ctx.Memory.TryWrite(errorAddress, error);
}
private static void WriteDecodeStreamResult(
CpuContext ctx,
ulong resultAddress,
int inputConsumed,
int outputWritten,
ulong totalDecodedSamples,
uint frames)
{
if (resultAddress == 0)
{
return;
}
Span<byte> sideband = stackalloc byte[DecodeSidebandBytes];
sideband.Clear();
// AjmSidebandResult.result / internal_result = 0 (OK)
BinaryPrimitives.WriteInt32LittleEndian(sideband.Slice(8, 4), inputConsumed);
BinaryPrimitives.WriteInt32LittleEndian(sideband.Slice(12, 4), outputWritten);
BinaryPrimitives.WriteUInt64LittleEndian(sideband.Slice(16, 8), totalDecodedSamples);
BinaryPrimitives.WriteUInt32LittleEndian(sideband.Slice(24, 4), frames);
_ = ctx.Memory.TryWrite(resultAddress, sideband);
}
private static void ClearGuestMemory(CpuContext ctx, ulong address, ulong byteCount)
{
if (address == 0 || byteCount == 0)
{
return;
}
var remaining = byteCount;
var cursor = address;
Span<byte> zero = stackalloc byte[256];
while (remaining > 0)
{
var chunk = (int)Math.Min(remaining, (ulong)zero.Length);
if (!ctx.Memory.TryWrite(cursor, zero[..chunk]))
{
return;
}
cursor += (ulong)chunk;
remaining -= (ulong)chunk;
}
}
private static ulong ReadStackArg64(CpuContext ctx, int index)
{
var address = ctx[CpuRegister.Rsp] + sizeof(ulong) + ((ulong)index * sizeof(ulong));
return TryReadUInt64(ctx, address, out var value) ? value : 0;
}
private static bool TryReadUInt64(CpuContext ctx, ulong address, out ulong value)
{
Span<byte> buffer = stackalloc byte[sizeof(ulong)];
if (!ctx.Memory.TryRead(address, buffer))
{
value = 0;
return false;
}
value = BinaryPrimitives.ReadUInt64LittleEndian(buffer);
return true;
}
private static bool TryWriteUInt64(CpuContext ctx, ulong address, ulong value)
{
Span<byte> buffer = stackalloc byte[sizeof(ulong)];
BinaryPrimitives.WriteUInt64LittleEndian(buffer, value);
return ctx.Memory.TryWrite(address, buffer);
}
private static void Trace(string message)
+14 -279
View File
@@ -4,7 +4,6 @@
using SharpEmu.HLE;
using SharpEmu.HLE.Host;
using System.Buffers;
using System.Buffers.Binary;
using System.Collections.Concurrent;
using System.Diagnostics;
@@ -12,17 +11,8 @@ namespace SharpEmu.Libs.Audio;
public static class AudioOutExports
{
private const int AudioOutOutputParamSize = 16;
private const int AudioOutMaximumOutputCount = 25;
internal const int AudioOutErrorInvalidPort = unchecked((int)0x80260003);
internal const int AudioOutErrorInvalidPointer = unchecked((int)0x80260004);
internal const int AudioOutErrorPortFull = unchecked((int)0x80260005);
internal const int AudioOutErrorInvalidSize = unchecked((int)0x80260006);
private static readonly ConcurrentDictionary<int, PortState> Ports = new();
private static int _nextPortHandle;
private static Func<uint, IHostAudioStream?>? _streamFactoryForTests;
// Diagnostic: confirm sceAudioOutOutput is actually called and whether the
// guest submits real samples or silence. Gated so it costs nothing when off.
@@ -66,7 +56,6 @@ public static class AudioOutExports
public int BytesPerSample { get; }
public bool IsFloat { get; }
public IHostAudioStream? Backend { get; }
public object SubmissionGate { get; } = new();
public volatile float Volume = 1.0f;
public int BufferByteLength =>
checked((int)BufferLength * Channels * BytesPerSample);
@@ -94,24 +83,7 @@ public static class AudioOutExports
}
}
public void Dispose()
{
lock (SubmissionGate)
{
Backend?.Dispose();
}
}
}
private readonly record struct OutputDescriptor(int Handle, ulong SourceAddress);
private struct ResolvedOutput
{
public int Handle;
public ulong SourceAddress;
public PortState Port;
public byte[]? HostBuffer;
public int HostBufferLength;
public void Dispose() => Backend?.Dispose();
}
[SysAbiExport(
@@ -143,18 +115,9 @@ public static class AudioOutExports
string backendName;
try
{
var streamFactory = Volatile.Read(ref _streamFactoryForTests);
if (streamFactory is not null)
{
backend = streamFactory(frequency);
backendName = "test";
}
else
{
var audio = HostPlatform.Current.Audio;
backend = audio.OpenStereoPcm16Stream(frequency);
backendName = audio.BackendName;
}
var audio = HostPlatform.Current.Audio;
backend = audio.OpenStereoPcm16Stream(frequency);
backendName = audio.BackendName;
}
catch (Exception exception)
{
@@ -229,46 +192,6 @@ public static class AudioOutExports
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "w3PdaSTSwGE",
ExportName = "sceAudioOutOutputs",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceAudioOut")]
public static int AudioOutOutputs(CpuContext ctx)
{
var parameterAddress = ctx[CpuRegister.Rdi];
var outputCount = unchecked((uint)ctx[CpuRegister.Rsi]);
if (outputCount == 0 || outputCount > AudioOutMaximumOutputCount)
{
return ctx.SetReturn(AudioOutErrorPortFull);
}
if (parameterAddress == 0)
{
return ctx.SetReturn(AudioOutErrorInvalidPointer);
}
var count = checked((int)outputCount);
Span<byte> parameterBytes =
stackalloc byte[AudioOutMaximumOutputCount * AudioOutOutputParamSize];
parameterBytes = parameterBytes[..checked(count * AudioOutOutputParamSize)];
if (!ctx.Memory.TryRead(parameterAddress, parameterBytes))
{
return ctx.SetReturn(AudioOutErrorInvalidPointer);
}
Span<OutputDescriptor> descriptors = stackalloc OutputDescriptor[count];
for (var i = 0; i < count; i++)
{
var entry = parameterBytes.Slice(i * AudioOutOutputParamSize, AudioOutOutputParamSize);
descriptors[i] = new OutputDescriptor(
BinaryPrimitives.ReadInt32LittleEndian(entry),
BinaryPrimitives.ReadUInt64LittleEndian(entry[8..]));
}
return ctx.SetReturn(SubmitOutputs(ctx, descriptors));
}
[SysAbiExport(
Nid = "QOQtbeDqsT4",
ExportName = "sceAudioOutOutput",
@@ -302,7 +225,16 @@ public static class AudioOutExports
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
TraceOutput(handle, port, source);
if (_traceOutput)
{
var n = Interlocked.Increment(ref _outputCount);
if (n <= 8 || n % 200 == 0)
{
var peak = PeakAmplitude(source, port.IsFloat, port.BytesPerSample);
Console.Error.WriteLine(
$"[LOADER][TRACE] audioout.output#{n} handle={handle} bytes={source.Length} ch={port.Channels} float={port.IsFloat} vol={port.Volume:F2} peak={peak:F4} backend={(port.Backend is null ? "none" : "coreaudio")}");
}
}
if (port.Backend is null)
{
@@ -340,184 +272,6 @@ public static class AudioOutExports
}
}
private static int SubmitOutputs(CpuContext ctx, ReadOnlySpan<OutputDescriptor> descriptors)
{
var resolvedArray = ArrayPool<ResolvedOutput>.Shared.Rent(descriptors.Length);
var resolved = resolvedArray.AsSpan(0, descriptors.Length);
resolved.Clear();
Span<int> lockOrder = stackalloc int[descriptors.Length];
var acquiredLocks = 0;
try
{
uint bufferLength = 0;
for (var i = 0; i < descriptors.Length; i++)
{
var descriptor = descriptors[i];
for (var previous = 0; previous < i; previous++)
{
if (resolved[previous].Handle == descriptor.Handle)
{
return AudioOutErrorInvalidPort;
}
}
if (!Ports.TryGetValue(descriptor.Handle, out var port))
{
return _shutdown ? 0 : AudioOutErrorInvalidPort;
}
if (i == 0)
{
bufferLength = port.BufferLength;
}
else if (port.BufferLength != bufferLength)
{
return AudioOutErrorInvalidSize;
}
resolved[i].Handle = descriptor.Handle;
resolved[i].SourceAddress = descriptor.SourceAddress;
resolved[i].Port = port;
lockOrder[i] = i;
}
// Every batch takes port locks in handle order. Two guest threads can
// submit overlapping batches in a different descriptor order without
// deadlocking each other.
for (var i = 1; i < lockOrder.Length; i++)
{
var index = lockOrder[i];
var position = i;
while (position > 0 &&
resolved[lockOrder[position - 1]].Handle > resolved[index].Handle)
{
lockOrder[position] = lockOrder[position - 1];
position--;
}
lockOrder[position] = index;
}
for (; acquiredLocks < lockOrder.Length; acquiredLocks++)
{
Monitor.Enter(resolved[lockOrder[acquiredLocks]].Port.SubmissionGate);
}
// AudioOutClose removes the handle before waiting for SubmissionGate.
// Recheck after acquiring all gates so a close racing this batch cannot
// turn a validated submission into a write to a disposed backend.
for (var i = 0; i < resolved.Length; i++)
{
if (!Ports.TryGetValue(resolved[i].Handle, out var current) ||
!ReferenceEquals(current, resolved[i].Port))
{
return _shutdown ? 0 : AudioOutErrorInvalidPort;
}
}
// Stage every guest buffer before the first host submission. A bad
// pointer in a later descriptor therefore cannot partially enqueue the
// earlier ports.
for (var i = 0; i < resolved.Length; i++)
{
ref var output = ref resolved[i];
if (output.SourceAddress == 0)
{
continue;
}
var sourceBuffer = ArrayPool<byte>.Shared.Rent(output.Port.BufferByteLength);
try
{
var source = sourceBuffer.AsSpan(0, output.Port.BufferByteLength);
if (!ctx.Memory.TryRead(output.SourceAddress, source))
{
return AudioOutErrorInvalidPointer;
}
TraceOutput(output.Handle, output.Port, source);
output.HostBufferLength = checked(
(int)output.Port.BufferLength * AudioPcmConversion.OutputFrameSize);
output.HostBuffer = ArrayPool<byte>.Shared.Rent(output.HostBufferLength);
AudioPcmConversion.ConvertToStereoPcm16(
source,
output.HostBuffer.AsSpan(0, output.HostBufferLength),
checked((int)output.Port.BufferLength),
output.Port.Channels,
output.Port.BytesPerSample,
output.Port.IsFloat,
output.Port.Volume);
}
finally
{
ArrayPool<byte>.Shared.Return(sourceBuffer);
}
}
PortState? pacingPort = null;
for (var i = 0; i < resolved.Length; i++)
{
ref var output = ref resolved[i];
if (output.HostBuffer is null ||
output.Port.Backend is null ||
!output.Port.Backend.Submit(
output.HostBuffer.AsSpan(0, output.HostBufferLength)))
{
if (pacingPort is null ||
HasLongerBufferDuration(output.Port, pacingPort))
{
pacingPort = output.Port;
}
}
}
// A batch is one guest scheduling point. When one or more ports have
// no usable backend, pace once using the longest affected buffer rather
// than sleeping once per port.
pacingPort?.PaceSilence();
return checked((int)resolved[0].Port.BufferLength);
}
finally
{
for (var i = acquiredLocks - 1; i >= 0; i--)
{
Monitor.Exit(resolved[lockOrder[i]].Port.SubmissionGate);
}
for (var i = 0; i < resolved.Length; i++)
{
if (resolved[i].HostBuffer is { } hostBuffer)
{
ArrayPool<byte>.Shared.Return(hostBuffer);
}
}
ArrayPool<ResolvedOutput>.Shared.Return(resolvedArray, clearArray: true);
}
}
private static bool HasLongerBufferDuration(PortState candidate, PortState current) =>
(ulong)candidate.BufferLength * current.Frequency >
(ulong)current.BufferLength * candidate.Frequency;
private static void TraceOutput(int handle, PortState port, ReadOnlySpan<byte> source)
{
if (!_traceOutput)
{
return;
}
var n = Interlocked.Increment(ref _outputCount);
if (n <= 8 || n % 200 == 0)
{
var peak = PeakAmplitude(source, port.IsFloat, port.BytesPerSample);
Console.Error.WriteLine(
$"[LOADER][TRACE] audioout.output#{n} handle={handle} bytes={source.Length} ch={port.Channels} float={port.IsFloat} vol={port.Volume:F2} peak={peak:F4} backend={(port.Backend is null ? "none" : "coreaudio")}");
}
}
[SysAbiExport(
Nid = "b+uAV89IlxE",
ExportName = "sceAudioOutSetVolume",
@@ -608,25 +362,6 @@ public static class AudioOutExports
}
}
internal static void SetStreamFactoryForTests(Func<uint, IHostAudioStream?>? streamFactory) =>
Volatile.Write(ref _streamFactoryForTests, streamFactory);
internal static void ResetForTests()
{
foreach (var handle in Ports.Keys)
{
if (Ports.TryRemove(handle, out var port))
{
port.Dispose();
}
}
_nextPortHandle = 0;
_outputCount = 0;
Volatile.Write(ref _shutdown, false);
Volatile.Write(ref _streamFactoryForTests, null);
}
private static bool _shutdown;
private static bool TryGetFormat(
+3 -16
View File
@@ -1,8 +1,6 @@
// 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
@@ -12,8 +10,7 @@ namespace SharpEmu.Libs.Gpu;
// translation for its API.
/// <summary>A guest texture referenced by a draw or dispatch. Format/NumberType/
/// TileMode/DstSelect/Type are raw guest descriptor codes. Depth is the
/// normalized volume depth (one for non-3D resources).</summary>
/// TileMode/DstSelect are raw guest descriptor codes.</summary>
internal sealed record GuestDrawTexture(
ulong Address,
uint Width,
@@ -35,15 +32,7 @@ internal sealed record GuestDrawTexture(
// from; -1 when the range is untracked or the pixels were not read here.
long WriteGeneration = -1,
bool ArrayedView = false,
uint ArrayLayers = 1,
uint Type = 9,
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);
uint ArrayLayers = 1);
/// <summary>Raw guest sampler descriptor dwords, copied verbatim from guest memory.</summary>
internal readonly record struct GuestSampler(
@@ -66,9 +55,7 @@ internal readonly record struct TextureContentIdentity(
uint Pitch,
GuestSampler Sampler,
bool Arrayed = false,
uint ArrayLayers = 1,
uint Type = 9,
uint Depth = 1);
uint ArrayLayers = 1);
internal sealed record GuestMemoryBuffer(
ulong BaseAddress,
@@ -1,313 +0,0 @@
// 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;
/// <summary>
/// Metal twin of <c>VulkanDetilePass</c>: runs the ExactXor detile equation from
/// <see cref="GnmTiling.GetDetileParams"/> as a Metal compute kernel
/// (<see cref="MslFixedShaders.CreateDetileCompute"/>), writing a linear buffer
/// and blitting it into the sampled texture.
///
/// <see cref="RecordDetile"/> 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 <b>Mac-untested</b> — mirrors the
/// verified Vulkan logic and the existing Metal message-send conventions, but
/// must be validated on a real Metal device.
/// </summary>
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;
/// <summary>
/// Records the deswizzle of <paramref name="tiled"/> into
/// <paramref name="texture"/> (<paramref name="texelWidth"/> x
/// <paramref name="texelHeight"/> texels x <paramref name="layers"/> slices)
/// onto <paramref name="commandBuffer"/>. The kernel iterates the element grid
/// from <paramref name="parameters"/> (for block-compressed formats a 4x4 block
/// is one element). Does not commit; the caller releases
/// <paramref name="transientBuffers"/> when the command buffer completes.
/// Returns false (empty transients) when unsupported or the pipeline could not
/// be built.
/// </summary>
public bool RecordDetile(
nint commandBuffer,
nint texture,
uint texelWidth,
uint texelHeight,
uint layers,
ReadOnlySpan<byte> 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<uint> 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;
}
}
}
@@ -1,7 +1,6 @@
// 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;
@@ -1584,45 +1583,6 @@ 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.
-12
View File
@@ -80,18 +80,6 @@ public static class JsonExports
return (int)OrbisGen2Result.ORBIS_GEN2_OK;
}
[SysAbiExport(
Nid = "PR5k1penBLM",
ExportName = "_ZN3sce4Json11Initializer9terminateEv",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceJson")]
public static int InitializerTerminate(CpuContext ctx)
{
TraceJson("Initializer.terminate", ctx[CpuRegister.Rdi], 0);
ctx[CpuRegister.Rax] = 0;
return (int)OrbisGen2Result.ORBIS_GEN2_OK;
}
[SysAbiExport(
Nid = "Cxwy7wHq4J0",
ExportName = "_ZN3sce4Json11Initializer10initializeEPKNS0_13InitParameterE",
@@ -153,64 +153,6 @@ public static class KernelPthreadCompatExports
static KernelPthreadCompatExports()
{
RunSynchronizationSelfChecks();
GuestThreadExecution.GuestThreadAbandoned += AbandonMutexesOwnedByThread;
}
/// <summary>
/// Force-release mutexes still owned by a guest thread that is being torn
/// down without a clean unlock (TBB worker_abort, abrupt exit). Otherwise
/// waiters can spin forever and block splash→first GPU submit.
/// </summary>
public static int AbandonMutexesOwnedByThread(ulong threadId, string reason)
{
if (threadId == 0)
{
return 0;
}
var released = 0;
var wakeKeys = new List<string>();
foreach (var pair in _mutexStates)
{
var state = pair.Value;
string? wakeKey = null;
lock (state)
{
if (state.OwnerThreadId != threadId || state.RecursionCount <= 0)
{
continue;
}
state.OwnerThreadId = 0;
state.RecursionCount = 0;
wakeKey = state.Waiters.First?.Value.Cooperative == true
? state.Waiters.First.Value.WakeKey
: null;
Monitor.PulseAll(state);
released++;
Console.Error.WriteLine(
$"[LOADER][WARN] pthread_mutex_abandon mutex=0x{pair.Key:X16} " +
$"owner={KernelPthreadState.DescribeThreadHandle(threadId)} " +
$"reason={reason} waiters={state.Waiters.Count}");
}
if (wakeKey is not null)
{
wakeKeys.Add(wakeKey);
}
}
foreach (var wakeKey in wakeKeys)
{
_ = GuestThreadExecution.Scheduler?.WakeBlockedThreads(wakeKey, 1);
}
if (released > 0)
{
Console.Error.Flush();
}
return released;
}
[SysAbiExport(
+3 -31
View File
@@ -27,12 +27,8 @@ internal static class KernelPthreadState
internal static ulong GetCurrentThreadHandle()
{
var guestThreadHandle = GuestThreadExecution.CurrentGuestThreadHandle;
// Prefer the bound guest handle even when it is not yet in Threads.
// Falling through to a synthetic ThreadStatic handle while a guest
// thread is bound causes mutex owner mismatches (unlock PERM → hang).
if (guestThreadHandle != 0)
if (guestThreadHandle != 0 && TryGetThreadIdentity(guestThreadHandle, out _))
{
EnsureGuestThreadIdentity(guestThreadHandle);
return guestThreadHandle;
}
@@ -43,27 +39,15 @@ internal static class KernelPthreadState
internal static ulong GetCurrentThreadUniqueId()
{
var guestThreadHandle = GuestThreadExecution.CurrentGuestThreadHandle;
if (guestThreadHandle != 0)
if (guestThreadHandle != 0 && TryGetThreadIdentity(guestThreadHandle, out var identity))
{
return EnsureGuestThreadIdentity(guestThreadHandle).UniqueId;
return identity.UniqueId;
}
EnsureCurrentThreadRegistered();
return _currentThreadUniqueId;
}
internal static string DescribeThreadHandle(ulong threadHandle)
{
if (threadHandle == 0)
{
return "none";
}
return TryGetThreadIdentity(threadHandle, out var identity)
? $"0x{threadHandle:X16}('{identity.Name}')"
: $"0x{threadHandle:X16}";
}
internal static ulong CreateThreadHandle(string name)
{
var uniqueId = unchecked((ulong)Interlocked.Increment(ref _nextUniqueThreadId));
@@ -75,18 +59,6 @@ internal static class KernelPthreadState
return Threads.TryGetValue(threadHandle, out identity);
}
private static ThreadIdentity EnsureGuestThreadIdentity(ulong guestThreadHandle)
{
if (Threads.TryGetValue(guestThreadHandle, out var existing))
{
return existing;
}
var uniqueId = unchecked((ulong)Interlocked.Increment(ref _nextUniqueThreadId));
var identity = new ThreadIdentity(uniqueId, $"Guest-0x{guestThreadHandle:X}");
return Threads.GetOrAdd(guestThreadHandle, identity);
}
private static void EnsureCurrentThreadRegistered()
{
if (_currentThreadHandle != 0)
-46
View File
@@ -755,52 +755,6 @@ public static class NetExports
return true;
}
[SysAbiExport(
Nid = "8Kcp5d-q1Uo",
ExportName = "sceNetInetPton",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceNet")]
public static int NetInetPton(CpuContext ctx)
{
var addressFamily = unchecked((int)ctx[CpuRegister.Rdi]);
var sourceAddress = ctx[CpuRegister.Rsi];
var destinationAddress = ctx[CpuRegister.Rdx];
if (sourceAddress == 0 || destinationAddress == 0)
{
return SetNetError(ctx, NetErrorInvalidArgument, NetErrnoInvalidArgument);
}
if (!TryReadUtf8Z(ctx, sourceAddress, MaxNameLength, out var source))
{
return SetNetError(ctx, NetErrorInvalidArgument, NetErrnoInvalidArgument);
}
var family = addressFamily switch
{
2 => AddressFamily.InterNetwork, // AF_INET
28 => AddressFamily.InterNetworkV6, // AF_INET6
_ => AddressFamily.Unknown,
};
if (family == AddressFamily.Unknown ||
!IPAddress.TryParse(source, out var parsed) ||
parsed.AddressFamily != family)
{
// Match BSD inet_pton: return 0 for a parseable-family miss.
ctx[CpuRegister.Rax] = 0;
return 0;
}
var bytes = parsed.GetAddressBytes();
if (!ctx.Memory.TryWrite(destinationAddress, bytes))
{
return SetNetError(ctx, NetErrorInvalidArgument, NetErrnoInvalidArgument);
}
TraceNet("inet_pton", addressFamily, sourceAddress, destinationAddress, (ulong)bytes.Length);
ctx[CpuRegister.Rax] = 1;
return 1;
}
private static void TraceNet(string operation, int id, ulong arg0, ulong arg1, ulong arg2)
{
if (!string.Equals(Environment.GetEnvironmentVariable("SHARPEMU_LOG_NET"), "1", StringComparison.Ordinal))
-365
View File
@@ -1,365 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using SharpEmu.HLE;
namespace SharpEmu.Libs.Psml;
public static class PsmlExports
{
// Empirically for Astro Bot (PPSA21567):
// [0] must be 0x80 (sizeThis / r9); any other value в†’ Allocate length=0
// AllocateMainDirectMemory(length=[0]*[16], alignment=[8], type=0xC)
// So put desired byte size in [8] (becomes alignment) and page size in [16].
private const ulong RequirementStructSize = 0x80;
private const ulong SharedResourcesPageSize = 0x10000;
private const ulong SharedResourcesBufferSizeBytes = 0x2000000;
private const ulong SharedResourcesContextSizeBytes = 0x100000;
private const ulong ContextBufferSizeBytes = 0x800000;
private const ulong ContextBufferAuxSizeBytes = 0x100000;
private static int _mfsrInitialized;
private static readonly Lock SharedResourcesGate = new();
private static readonly Dictionary<ulong, SharedResourcesState> SharedResourcesByDescriptor = new();
private static readonly Lock ContextGate = new();
private static readonly Dictionary<ulong, ContextState> ContextsByAddress = new();
[SysAbiExport(
Nid = "3WVD91e12ZQ",
ExportName = "scePsmlMfsrInit",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrInit(CpuContext ctx)
{
var arg0 = ctx[CpuRegister.Rdi];
var arg1 = ctx[CpuRegister.Rsi];
var arg2 = ctx[CpuRegister.Rdx];
Interlocked.Exchange(ref _mfsrInitialized, 1);
TracePsml($"mfsr_init arg0=0x{arg0:X} arg1=0x{arg1:X} arg2=0x{arg2:X}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "+2KpvixvL6E",
ExportName = "scePsmlMfsrGetSharedResourcesInitRequirement",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetSharedResourcesInitRequirement(CpuContext ctx)
{
var bufferRequirementAddress = ctx[CpuRegister.Rdi];
var contextRequirementAddress = ctx[CpuRegister.Rsi];
var flags = ctx[CpuRegister.Rdx];
var configAddress = ctx[CpuRegister.Rcx];
if (bufferRequirementAddress == 0 || contextRequirementAddress == 0)
{
return ctx.SetReturn(OrbisGen2Result.ORBIS_GEN2_ERROR_INVALID_ARGUMENT);
}
if (!WriteMemoryRequirement(ctx, bufferRequirementAddress, SharedResourcesBufferSizeBytes) ||
!WriteMemoryRequirement(ctx, contextRequirementAddress, SharedResourcesContextSizeBytes))
{
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
TracePsml(
$"mfsr_get_shared_resources_init_requirement buf=0x{bufferRequirementAddress:X16} " +
$"ctx=0x{contextRequirementAddress:X16} flags=0x{flags:X} config=0x{configAddress:X16} " +
$"buf_size=0x{SharedResourcesBufferSizeBytes:X} ctx_size=0x{SharedResourcesContextSizeBytes:X}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "eWoKNeB6V-k",
ExportName = "scePsmlMfsrCreateSharedResources",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrCreateSharedResources(CpuContext ctx)
{
var descriptorAddress = ctx[CpuRegister.Rdi];
var contextRequirementAddress = ctx[CpuRegister.Rsi];
var directMemoryAddress = ctx[CpuRegister.Rdx];
if (descriptorAddress == 0 || contextRequirementAddress == 0 || directMemoryAddress == 0)
{
return ctx.SetReturn(OrbisGen2Result.ORBIS_GEN2_ERROR_INVALID_ARGUMENT);
}
var contextSizeBytes = ReadRequirementSize(ctx, contextRequirementAddress, SharedResourcesContextSizeBytes);
var state = new SharedResourcesState(
DescriptorAddress: descriptorAddress,
DirectMemoryAddress: directMemoryAddress,
BufferSizeBytes: SharedResourcesBufferSizeBytes,
ContextSizeBytes: contextSizeBytes,
PageSizeBytes: SharedResourcesPageSize);
if (!WriteSharedResourcesDescriptor(ctx, state))
{
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
lock (SharedResourcesGate)
{
SharedResourcesByDescriptor[descriptorAddress] = state;
}
TracePsml(
$"mfsr_create_shared_resources desc=0x{descriptorAddress:X16} req=0x{contextRequirementAddress:X16} " +
$"direct=0x{directMemoryAddress:X16} buf_size=0x{state.BufferSizeBytes:X} " +
$"ctx_size=0x{state.ContextSizeBytes:X} page=0x{state.PageSizeBytes:X}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "ArakEpzsZo0",
ExportName = "scePsmlMfsrGetContextBufferRequirement800M3_2",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetContextBufferRequirement800M3_2(CpuContext ctx)
{
var bufferRequirementAddress = ctx[CpuRegister.Rdi];
var contextRequirementAddress = ctx[CpuRegister.Rsi];
var directMemoryAddress = ctx[CpuRegister.Rdx];
if (bufferRequirementAddress == 0 || contextRequirementAddress == 0 || directMemoryAddress == 0)
{
return ctx.SetReturn(OrbisGen2Result.ORBIS_GEN2_ERROR_INVALID_ARGUMENT);
}
if (!WriteMemoryRequirement(ctx, bufferRequirementAddress, ContextBufferSizeBytes) ||
!WriteMemoryRequirement(ctx, contextRequirementAddress, ContextBufferAuxSizeBytes))
{
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
TracePsml(
$"mfsr_get_context_buffer_requirement_800m3_2 buf=0x{bufferRequirementAddress:X16} " +
$"ctx=0x{contextRequirementAddress:X16} direct=0x{directMemoryAddress:X16} " +
$"buf_size=0x{ContextBufferSizeBytes:X} aux_size=0x{ContextBufferAuxSizeBytes:X}");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "gxv3i+MTEzU",
ExportName = "scePsmlMfsrCreateContext800M3_2",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrCreateContext800M3_2(CpuContext ctx)
{
var contextAddress = ctx[CpuRegister.Rdi];
var requirementAddress = ctx[CpuRegister.Rsi];
var structSize = ctx[CpuRegister.Rdx];
var sharedDirectMemory = ctx[CpuRegister.Rcx];
var pageSize = ctx[CpuRegister.R8];
if (contextAddress == 0 || sharedDirectMemory == 0)
{
return ctx.SetReturn(OrbisGen2Result.ORBIS_GEN2_ERROR_INVALID_ARGUMENT);
}
var sharedState = TryFindSharedResources(sharedDirectMemory, contextAddress);
var effectiveStructSize = structSize != 0 ? structSize : RequirementStructSize;
var effectivePageSize = pageSize != 0 ? pageSize : SharedResourcesPageSize;
var sharedDescriptor = sharedState?.DescriptorAddress ?? contextAddress - 0x30;
var bufferSize = sharedState?.BufferSizeBytes ?? ContextBufferSizeBytes;
var state = new ContextState(
ContextAddress: contextAddress,
SharedResourcesDescriptor: sharedDescriptor,
DirectMemoryAddress: sharedDirectMemory,
BufferSizeBytes: bufferSize,
PageSizeBytes: effectivePageSize,
StructSizeBytes: effectiveStructSize);
if (!WriteContextObject(ctx, state))
{
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
lock (ContextGate)
{
ContextsByAddress[contextAddress] = state;
}
TracePsml(
$"mfsr_create_context_800m3_2 ctx=0x{contextAddress:X16} req=0x{requirementAddress:X16} " +
$"struct=0x{effectiveStructSize:X} direct=0x{sharedDirectMemory:X16} " +
$"shared_desc=0x{sharedDescriptor:X16} buf_size=0x{bufferSize:X} page=0x{effectivePageSize:X}");
return ctx.SetReturn(0);
}
private static bool WriteMemoryRequirement(CpuContext ctx, ulong address, ulong sizeBytes)
{
return ctx.TryWriteUInt64(address, RequirementStructSize) &&
ctx.TryWriteUInt64(address + 0x08, sizeBytes) &&
ctx.TryWriteUInt64(address + 0x10, SharedResourcesPageSize);
}
private static ulong ReadRequirementSize(CpuContext ctx, ulong address, ulong fallback)
{
if (!ctx.TryReadUInt64(address + 0x08, out var sizeBytes) || sizeBytes == 0)
{
return fallback;
}
// Guest requirement structs use size @8; reject pointer-like garbage.
return sizeBytes > 0x1000_0000UL ? fallback : sizeBytes;
}
private static SharedResourcesState? TryFindSharedResources(ulong directMemoryAddress, ulong contextAddress)
{
lock (SharedResourcesGate)
{
foreach (var state in SharedResourcesByDescriptor.Values)
{
if (state.DirectMemoryAddress == directMemoryAddress)
{
return state;
}
}
var inferredDescriptor = contextAddress >= 0x30 ? contextAddress - 0x30 : 0;
if (inferredDescriptor != 0 &&
SharedResourcesByDescriptor.TryGetValue(inferredDescriptor, out var byDescriptor))
{
return byDescriptor;
}
}
return null;
}
private static bool WriteContextObject(CpuContext ctx, ContextState state)
{
return ctx.TryWriteUInt64(state.ContextAddress + 0x00, state.StructSizeBytes) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x08, state.DirectMemoryAddress) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x10, state.SharedResourcesDescriptor) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x18, state.BufferSizeBytes) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x20, state.PageSizeBytes) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x28, state.ContextAddress) &&
ctx.TryWriteUInt64(state.ContextAddress + 0x30, state.DirectMemoryAddress);
}
private static bool WriteSharedResourcesDescriptor(CpuContext ctx, SharedResourcesState state)
{
// Stamp a compact self-describing blob so follow-up PSML calls can treat the
// descriptor as initialized guest memory instead of an all-zero placeholder.
return ctx.TryWriteUInt64(state.DescriptorAddress + 0x00, RequirementStructSize) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x08, state.DirectMemoryAddress) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x10, state.DirectMemoryAddress) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x18, state.BufferSizeBytes) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x20, state.ContextSizeBytes) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x28, state.PageSizeBytes) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x30, state.DirectMemoryAddress) &&
ctx.TryWriteUInt64(state.DescriptorAddress + 0x38, state.BufferSizeBytes + state.ContextSizeBytes);
}
// Astro logo path: SizeInDwords then GetDispatchMfsrPacket900 (RUNLFro+qok).
// Unresolved 900 returns non-zero and trips GfxRenderStagePSSR.cpp:266.
private const int SoftPacketSizeInDwords = 0x80;
private const ulong SoftPacketSizeBytes = SoftPacketSizeInDwords * 4UL;
[SysAbiExport(
Nid = "AHalTX9wFZY",
ExportName = "scePsmlMfsrGetDispatchMfsrPacketSizeInDwords",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetDispatchMfsrPacketSizeInDwords(CpuContext ctx)
{
// Logo/PSSR path: guest asserts ret == 0, then calls GetDispatchMfsrPacket900
// with rdi=SoftPacketSizeInDwords (observed 0x80). Returning the size in rax
// tripped :266 and skipped the 900 call (tDISP-s6). SCE_OK keeps the chain.
var arg0 = ctx[CpuRegister.Rdi];
TracePsml(
$"mfsr_get_dispatch_packet_size_dwords arg0=0x{arg0:X} " +
$"size=0x{SoftPacketSizeInDwords:X} ret=0");
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "RUNLFro+qok",
ExportName = "scePsmlMfsrGetDispatchMfsrPacket900",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetDispatchMfsrPacket900(CpuContext ctx) =>
SoftGetDispatchMfsrPacket(ctx, "900");
[SysAbiExport(
Nid = "s2psNHUIdjk",
ExportName = "scePsmlMfsrGetDispatchMfsrPacket1000",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetDispatchMfsrPacket1000(CpuContext ctx) =>
SoftGetDispatchMfsrPacket(ctx, "1000");
[SysAbiExport(
Nid = "94iBp3KvIuI",
ExportName = "scePsmlMfsrGetDispatchMfsrPacket1100",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libScePsml")]
public static int PsmlMfsrGetDispatchMfsrPacket1100(CpuContext ctx) =>
SoftGetDispatchMfsrPacket(ctx, "1100");
private static int SoftGetDispatchMfsrPacket(CpuContext ctx, string variant)
{
var arg0 = ctx[CpuRegister.Rdi];
var arg1 = ctx[CpuRegister.Rsi];
var arg2 = ctx[CpuRegister.Rdx];
var arg3 = ctx[CpuRegister.Rcx];
// Logo call shape (tDISP-s3): rdi=size_dwords (0x80), rsi/rdx = guest
// packet/param buffers. Clear the first mapped buffer arg.
var packetAddress = 0UL;
foreach (var candidate in new[] { arg1, arg2, arg3 })
{
if (candidate >= 0x10000)
{
packetAddress = candidate;
break;
}
}
var cleared = packetAddress != 0 && TryClearGuestBuffer(ctx, packetAddress, SoftPacketSizeBytes);
TracePsml(
$"mfsr_get_dispatch_packet_{variant} a0=0x{arg0:X16} a1=0x{arg1:X16} " +
$"a2=0x{arg2:X16} a3=0x{arg3:X16} packet=0x{packetAddress:X16} cleared={cleared}");
return ctx.SetReturn(0);
}
private static bool TryClearGuestBuffer(CpuContext ctx, ulong address, ulong length)
{
Span<byte> zeroes = stackalloc byte[4096];
zeroes.Clear();
for (ulong offset = 0; offset < length;)
{
var chunkSize = (int)Math.Min((ulong)zeroes.Length, length - offset);
if (!ctx.Memory.TryWrite(address + offset, zeroes[..chunkSize]))
{
return false;
}
offset += unchecked((uint)chunkSize);
}
return true;
}
private static void TracePsml(string message)
{
if (string.Equals(Environment.GetEnvironmentVariable("SHARPEMU_LOG_PSML"), "1", StringComparison.Ordinal))
{
Console.Error.WriteLine($"[LOADER][TRACE] psml.{message}");
}
}
private readonly record struct SharedResourcesState(
ulong DescriptorAddress,
ulong DirectMemoryAddress,
ulong BufferSizeBytes,
ulong ContextSizeBytes,
ulong PageSizeBytes);
private readonly record struct ContextState(
ulong ContextAddress,
ulong SharedResourcesDescriptor,
ulong DirectMemoryAddress,
ulong BufferSizeBytes,
ulong PageSizeBytes,
ulong StructSizeBytes);
}
@@ -59,65 +59,4 @@ public static class GameServiceStubs
public static int NpUniversalDataSystemCreateEvent(CpuContext ctx) => OkWithHandle(ctx, CpuRegister.Rdi);
public static int NpUniversalDataSystemPostEvent(CpuContext ctx) => Ok(ctx);
public static int NpUniversalDataSystemDestroyEvent(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "47UAEuQl+iI", ExportName = "sceNpUniversalDataSystemTerminate",
Target = Generation.Gen5, LibraryName = "libSceNpUniversalDataSystem")]
public static int NpUniversalDataSystemTerminate(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "0HBYxYAjmf0", ExportName = "sceNpGameIntentTerminate",
Target = Generation.Gen5, LibraryName = "libSceNpGameIntent")]
public static int NpGameIntentTerminate(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "jqb7HntFQFc", ExportName = "sceWebBrowserDialogInitialize",
Target = Generation.Gen5, LibraryName = "libSceWebBrowserDialog")]
public static int WebBrowserDialogInitialize(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "ocHtyBwHfys", ExportName = "sceWebBrowserDialogTerminate",
Target = Generation.Gen5, LibraryName = "libSceWebBrowserDialog")]
public static int WebBrowserDialogTerminate(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "kvYEw2lBndk", ExportName = "sceGameLiveStreamingInitialize",
Target = Generation.Gen5, LibraryName = "libSceGameLiveStreaming")]
public static int GameLiveStreamingInitialize(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "isruqthpYcw", ExportName = "sceSharePlayInitialize",
Target = Generation.Gen5, LibraryName = "libSceSharePlay")]
public static int SharePlayInitialize(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "0IL1keINExQ", ExportName = "sceShareTerminate",
Target = Generation.Gen5, LibraryName = "libSceShareUtility")]
public static int ShareTerminate(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "YBiIdcDPrxs", ExportName = "sceShareFeaturePermit",
Target = Generation.Gen5, LibraryName = "libSceShareUtility")]
public static int ShareFeaturePermit(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "9TrhuGzberQ", ExportName = "sceVoiceInit",
Target = Generation.Gen5, LibraryName = "libSceVoice")]
public static int VoiceInit(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "clyKUyi3RYU", ExportName = "sceVoiceSetThreadsParams",
Target = Generation.Gen5, LibraryName = "libSceVoice")]
public static int VoiceSetThreadsParams(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "dPj4ZtRcIWk", ExportName = "sceContentSearchInit",
Target = Generation.Gen5, LibraryName = "libSceContentSearch")]
public static int ContentSearchInit(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "zoxb0wEChEM", ExportName = "sceContentDeleteInitialize",
Target = Generation.Gen5, LibraryName = "libSceContentDelete")]
public static int ContentDeleteInitialize(CpuContext ctx) => Ok(ctx);
[SysAbiExport(Nid = "Fc8qxlKINYQ", ExportName = "sceVideoRecordingSetInfo",
Target = Generation.Gen5, LibraryName = "libSceVideoRecording")]
public static int VideoRecordingSetInfo(CpuContext ctx) => Ok(ctx);
// Captured from GTA V Enhanced (PPSA04264); not in the public NID catalog.
// Side-effect-free success — same as unresolved stub behavior that kept boot
// moving; reverse the ABI before writing guest memory.
#pragma warning disable SHEM006
[SysAbiExport(Nid = "Ikfdt-rIqCE", ExportName = "sceUnknownIkfdt",
Target = Generation.Gen5, LibraryName = "libKernel")]
public static int UnknownIkfdt(CpuContext ctx) => Ok(ctx);
#pragma warning restore SHEM006
}
@@ -18,7 +18,6 @@ public static class SystemServiceExports
private const int TitleIdFieldSize = 0x10;
private static string? _mainAppTitleId;
private static int _noticeScreenSkipFlag;
public static void ConfigureApplicationInfo(string? titleId)
{
@@ -38,40 +37,14 @@ public static class SystemServiceExports
return ctx.SetReturn(OrbisSystemServiceErrorParameter);
}
// Keep the flag state even though the emulator does not display the
// system notice screen. Titles use this service as a normal preference
// store and expect a later get to observe the value they set.
// No system notice screen to skip in the emulator; report "do not skip".
Span<byte> flagBytes = stackalloc byte[1];
flagBytes[0] = unchecked((byte)Volatile.Read(ref _noticeScreenSkipFlag));
flagBytes[0] = 0;
return ctx.Memory.TryWrite(flagAddress, flagBytes)
? ctx.SetReturn(0)
: ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
[SysAbiExport(
Nid = "8Lo6Zv94aho",
ExportName = "sceSystemServiceDisableNoticeScreenSkipFlagAutoSet",
Target = Generation.Gen5,
LibraryName = "libSceSystemService")]
public static int SystemServiceDisableNoticeScreenSkipFlagAutoSet(CpuContext ctx) =>
ctx.SetReturn(0);
// Settings entry calls this immediately before spawning SaveModTime/Load
// threads. An unresolved stub returns NOT_FOUND and the title can stall in
// that path; accept the write and report success.
[SysAbiExport(
Nid = "Q3utJvma4Mo",
ExportName = "sceSystemServiceSetNoticeScreenSkipFlag",
Target = Generation.Gen5,
LibraryName = "libSceSystemService")]
public static int SystemServiceSetNoticeScreenSkipFlag(CpuContext ctx)
{
// The native API takes the flag value in the first argument. Treat any
// non-zero value as true, matching the bool-like PS5 ABI.
Volatile.Write(ref _noticeScreenSkipFlag, ctx[CpuRegister.Rdi] != 0 ? 1 : 0);
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "4veE0XiIugA",
ExportName = "sceSystemServiceGetMainAppTitleId",
@@ -239,7 +212,4 @@ public static class SystemServiceExports
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceSystemService")]
public static int SystemServiceReportAbnormalTermination(CpuContext ctx) => ctx.SetReturn(0);
internal static void ResetForTests() =>
Volatile.Write(ref _noticeScreenSkipFlag, 0);
}
@@ -1,894 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Numerics;
using System.Runtime.CompilerServices;
using SharpEmu.Libs.Agc;
using SharpEmu.ShaderCompiler.Vulkan;
using Silk.NET.Vulkan;
using VkBuffer = Silk.NET.Vulkan.Buffer;
namespace SharpEmu.Libs.VideoOut;
/// <summary>
/// Self-contained GPU deswizzle pass: runs the ExactXor detile equation from
/// <see cref="GnmTiling.GetDetileParams"/> as a Vulkan compute shader
/// (<see cref="SpirvFixedShaders.CreateDetileCompute"/>), writing a linear buffer
/// and copying it into a sampled image — the GPU equivalent of the CPU
/// <c>GnmTiling.TryDetile</c> + staging upload.
///
/// Two entry points share the same (verified) recording:
/// <see cref="DetileIntoImage"/> is a self-contained one-shot (submit + wait) used
/// by the isolation self-test; <see cref="RecordDetile"/> 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; <see cref="Supports"/> lets
/// the caller fall back to the CPU path for everything else.
/// </summary>
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;
private PhysicalDeviceMemoryProperties _memoryProperties;
private bool _memoryPropertiesLoaded;
private readonly Dictionary<int[], TermBuffer> _xorTermBuffers = new(ReferenceComparer.Instance);
private readonly Dictionary<int[], TermBuffer> _blockTermBuffers = new(ReferenceComparer.Instance);
private TermBuffer _placeholderTermBuffer;
private readonly Dictionary<(ulong Bucket, bool HostVisible), Stack<Allocation>> _bufferPool = new();
private readonly List<Allocation> _allAllocations = new();
private readonly Stack<DescriptorSet> _freeDescriptorSets = new();
private readonly List<DescriptorPool> _descriptorPools = new();
private const uint DescriptorSetsPerPool = 64;
private const ulong MinimumBufferBucket = 4096;
public VulkanDetilePass(
Vk vk,
Device device,
Queue queue,
PhysicalDevice physicalDevice,
uint queueFamilyIndex)
{
_vk = vk;
_device = device;
_queue = queue;
_physicalDevice = physicalDevice;
_queueFamilyIndex = queueFamilyIndex;
}
/// <summary>
/// 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.
/// </summary>
public static bool Supports(in DetileParams parameters) =>
(parameters.Equation == DetileEquation.ExactXor ||
parameters.Equation == DetileEquation.BlockTable) &&
parameters.BytesPerElement is 4 or 8 or 16;
/// <summary>Opaque handle to the pooled resources one recorded detile is using.
/// The caller hands it back to <see cref="Retire"/> once the command buffer they
/// were recorded into has completed; nothing is destroyed, the buffers and the
/// descriptor set return to this pass's free lists for the next texture.</summary>
public sealed class Transients
{
internal static readonly Transients Empty = new();
internal Allocation Tiled;
internal Allocation Output;
internal DescriptorSet Set;
internal bool Rented;
}
internal readonly record struct Allocation(
VkBuffer Buffer,
DeviceMemory Memory,
ulong Capacity,
bool HostVisible);
private readonly record struct TermBuffer(VkBuffer Buffer, ulong ByteSize);
private struct DetileResources
{
public Allocation Tiled;
public Allocation Output;
public DescriptorSet Set;
public ulong OutputBytes;
public uint SrcSliceElements;
public uint EquationValue;
public uint UintsPerElement;
}
private sealed class ReferenceComparer : IEqualityComparer<int[]>
{
public static readonly ReferenceComparer Instance = new();
public bool Equals(int[]? x, int[]? y) => ReferenceEquals(x, y);
public int GetHashCode(int[] obj) => RuntimeHelpers.GetHashCode(obj);
}
/// <summary>
/// Records the deswizzle of <paramref name="tiled"/> into <paramref name="image"/>
/// (<paramref name="texelWidth"/> x <paramref name="texelHeight"/> texels x
/// <paramref name="layers"/> array slices, currently in
/// <paramref name="currentLayout"/>) onto <paramref name="commandBuffer"/>,
/// leaving the image <see cref="ImageLayout.ShaderReadOnlyOptimal"/>. The kernel
/// iterates the element grid from <paramref name="parameters"/> (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 <paramref name="transients"/> with the command buffer's fence.
/// Returns false (with empty transients) when unsupported.
/// </summary>
public bool RecordDetile(
CommandBuffer commandBuffer,
Image image,
ImageLayout currentLayout,
uint texelWidth,
uint texelHeight,
uint layers,
ReadOnlySpan<byte> tiled,
in DetileParams parameters,
out Transients transients)
{
transients = Transients.Empty;
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
{
ReleaseResources(in resources);
throw;
}
transients = new Transients
{
Tiled = resources.Tiled,
Output = resources.Output,
Set = resources.Set,
Rented = true,
};
return true;
}
public void Retire(Transients transients)
{
if (_disposed || transients is null || !transients.Rented)
{
return;
}
transients.Rented = false;
ReturnBuffer(transients.Tiled);
ReturnBuffer(transients.Output);
_freeDescriptorSets.Push(transients.Set);
}
/// <summary>
/// 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 <see cref="RecordDetile"/> there.
/// </summary>
public bool DetileIntoImage(
Image image,
ImageLayout currentLayout,
uint texelWidth,
uint texelHeight,
uint layers,
ReadOnlySpan<byte> 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);
}
ReleaseResources(in resources);
}
}
private void PrepareResources(ReadOnlySpan<byte> 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.
TermBuffer xTerm;
TermBuffer yTerm;
if (parameters.Equation == DetileEquation.BlockTable)
{
xTerm = GetTermBuffer(_blockTermBuffers, parameters.BlockTable, shift: 0);
yTerm = GetPlaceholderTermBuffer();
resources.EquationValue = 1;
}
else
{
var shift = BitOperations.TrailingZeroCount((uint)parameters.BytesPerElement);
xTerm = GetTermBuffer(_xorTermBuffers, parameters.XByteTerm, shift);
yTerm = GetTermBuffer(_xorTermBuffers, 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 = RentBuffer((ulong)tiled.Length, hostVisible: true);
UploadBytes(resources.Tiled.Memory, tiled);
resources.Output = RentBuffer(resources.OutputBytes, hostVisible: false);
resources.Set = RentDescriptorSet();
WriteDescriptors(
resources.Set,
(resources.Tiled.Buffer, (ulong)tiled.Length),
(xTerm.Buffer, xTerm.ByteSize),
(yTerm.Buffer, yTerm.ByteSize),
(resources.Output.Buffer, resources.OutputBytes));
}
private TermBuffer GetTermBuffer(Dictionary<int[], TermBuffer> cache, int[] table, int shift)
{
if (cache.TryGetValue(table, out var cached))
{
return cached;
}
var terms = ToElementTerms(table, shift);
var byteSize = (ulong)terms.Length * sizeof(uint);
var allocation = CreateBuffer(byteSize, hostVisible: true);
UploadUInts(allocation.Memory, terms);
var termBuffer = new TermBuffer(allocation.Buffer, byteSize);
cache[table] = termBuffer;
return termBuffer;
}
private TermBuffer GetPlaceholderTermBuffer()
{
if (_placeholderTermBuffer.Buffer.Handle != 0)
{
return _placeholderTermBuffer;
}
var allocation = CreateBuffer(sizeof(uint), hostVisible: true);
UploadUInts(allocation.Memory, [0u]);
_placeholderTermBuffer = new TermBuffer(allocation.Buffer, sizeof(uint));
return _placeholderTermBuffer;
}
private static ulong BucketFor(ulong size)
{
var bucket = MinimumBufferBucket;
while (bucket < size)
{
bucket <<= 1;
}
return bucket;
}
private Allocation RentBuffer(ulong size, bool hostVisible)
{
var bucket = BucketFor(size);
if (_bufferPool.TryGetValue((bucket, hostVisible), out var free) && free.Count > 0)
{
return free.Pop();
}
return CreateBuffer(bucket, hostVisible);
}
private void ReturnBuffer(Allocation allocation)
{
if (allocation.Buffer.Handle == 0)
{
return;
}
var key = (allocation.Capacity, allocation.HostVisible);
if (!_bufferPool.TryGetValue(key, out var free))
{
free = new Stack<Allocation>();
_bufferPool[key] = free;
}
free.Push(allocation);
}
private DescriptorSet RentDescriptorSet()
{
if (_freeDescriptorSets.Count > 0)
{
return _freeDescriptorSets.Pop();
}
var poolSize = new DescriptorPoolSize
{
Type = DescriptorType.StorageBuffer,
DescriptorCount = 4 * DescriptorSetsPerPool,
};
var poolInfo = new DescriptorPoolCreateInfo
{
SType = StructureType.DescriptorPoolCreateInfo,
MaxSets = DescriptorSetsPerPool,
PoolSizeCount = 1,
PPoolSizes = &poolSize,
};
Check(
_vk.CreateDescriptorPool(_device, &poolInfo, null, out var pool),
"vkCreateDescriptorPool(detile)");
_descriptorPools.Add(pool);
var layouts = stackalloc DescriptorSetLayout[(int)DescriptorSetsPerPool];
for (var index = 0; index < DescriptorSetsPerPool; index++)
{
layouts[index] = _descriptorSetLayout;
}
var sets = stackalloc DescriptorSet[(int)DescriptorSetsPerPool];
var allocateInfo = new DescriptorSetAllocateInfo
{
SType = StructureType.DescriptorSetAllocateInfo,
DescriptorPool = pool,
DescriptorSetCount = DescriptorSetsPerPool,
PSetLayouts = layouts,
};
Check(
_vk.AllocateDescriptorSets(_device, &allocateInfo, sets),
"vkAllocateDescriptorSets(detile)");
for (var index = 0; index < DescriptorSetsPerPool; index++)
{
_freeDescriptorSets.Push(sets[index]);
}
return _freeDescriptorSets.Pop();
}
private void ReleaseResources(in DetileResources resources)
{
ReturnBuffer(resources.Tiled);
ReturnBuffer(resources.Output);
if (resources.Set.Handle != 0)
{
_freeDescriptorSets.Push(resources.Set);
}
}
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<uint> 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.Buffer,
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.Buffer, image, ImageLayout.TransferDstOptimal, 1, &copyRegion);
TransitionImage(
commandBuffer,
image,
ImageLayout.TransferDstOptimal,
ImageLayout.ShaderReadOnlyOptimal,
AccessFlags.TransferWriteBit,
AccessFlags.ShaderReadBit,
PipelineStageFlags.TransferBit,
PipelineStageFlags.FragmentShaderBit,
layers);
}
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<byte> 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 Allocation CreateBuffer(ulong size, bool hostVisible)
{
var bufferInfo = new BufferCreateInfo
{
SType = StructureType.BufferCreateInfo,
Size = size,
Usage = BufferUsageFlags.StorageBufferBit | BufferUsageFlags.TransferSrcBit,
SharingMode = SharingMode.Exclusive,
};
Check(_vk.CreateBuffer(_device, &bufferInfo, null, out var buffer), "vkCreateBuffer(detile)");
_vk.GetBufferMemoryRequirements(_device, buffer, out var requirements);
var required = hostVisible
? MemoryPropertyFlags.HostVisibleBit | MemoryPropertyFlags.HostCoherentBit
: MemoryPropertyFlags.DeviceLocalBit;
var allocateInfo = new MemoryAllocateInfo
{
SType = StructureType.MemoryAllocateInfo,
AllocationSize = requirements.Size,
MemoryTypeIndex = FindMemoryType(requirements.MemoryTypeBits, required, hostVisible),
};
Check(_vk.AllocateMemory(_device, &allocateInfo, null, out var memory), "vkAllocateMemory(detile)");
Check(_vk.BindBufferMemory(_device, buffer, memory, 0), "vkBindBufferMemory(detile)");
var allocation = new Allocation(buffer, memory, size, hostVisible);
_allAllocations.Add(allocation);
return allocation;
}
private uint FindMemoryType(uint typeBits, MemoryPropertyFlags requiredFlags, bool hostVisible)
{
if (!_memoryPropertiesLoaded)
{
_vk.GetPhysicalDeviceMemoryProperties(_physicalDevice, out _memoryProperties);
_memoryPropertiesLoaded = true;
}
fixed (PhysicalDeviceMemoryProperties* properties = &_memoryProperties)
{
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;
}
}
if (!hostVisible)
{
for (uint index = 0; index < properties->MemoryTypeCount; index++)
{
if ((typeBits & (1u << (int)index)) != 0)
{
return index;
}
}
}
}
throw new InvalidOperationException("No compatible Vulkan memory type for detile.");
}
private void UploadBytes(DeviceMemory memory, ReadOnlySpan<byte> data)
{
void* mapped;
Check(_vk.MapMemory(_device, memory, 0, (ulong)data.Length, 0, &mapped), "vkMapMemory(detile)");
data.CopyTo(new Span<byte>(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<uint>(mapped, data.Length));
_vk.UnmapMemory(_device, memory);
}
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;
foreach (var allocation in _allAllocations)
{
DestroyBuffer(allocation.Buffer, allocation.Memory);
}
_allAllocations.Clear();
_bufferPool.Clear();
_xorTermBuffers.Clear();
_blockTermBuffers.Clear();
_placeholderTermBuffer = default;
_freeDescriptorSets.Clear();
foreach (var pool in _descriptorPools)
{
_vk.DestroyDescriptorPool(_device, pool, null);
}
_descriptorPools.Clear();
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);
}
}
}
@@ -1,482 +0,0 @@
// 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;
/// <summary>
/// Opt-in GPU equivalence check for <see cref="VulkanDetilePass"/>. 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
/// <see cref="GnmTiling.TryDetile"/> — 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.
/// </summary>
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<Image, bool> 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<byte>(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<byte> 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);
pass.Retire(transients);
}
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,
&region);
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}");
}
}
}
File diff suppressed because it is too large Load Diff
@@ -16,37 +16,4 @@ public static class VoiceQoSExports
{
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "Trpt2QBZHCI",
ExportName = "sceVoiceQoSGetStatus",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceVoiceQoS")]
public static int VoiceQoSGetStatus(CpuContext ctx)
{
// Returns 0 to indicate connected state (voice available)
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "FuXenJLkk-c",
ExportName = "sceVoiceQoSTerminate",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceVoiceQoS")]
public static int VoiceQoSTerminate(CpuContext ctx)
{
// No-op: cleanup is handled by emulator shutdown
return ctx.SetReturn(0);
}
[SysAbiExport(
Nid = "+0lOiPZjnBI",
ExportName = "sceVoiceQoSSetMode",
Target = Generation.Gen4 | Generation.Gen5,
LibraryName = "libSceVoiceQoS")]
public static int VoiceQoSSetMode(CpuContext ctx)
{
// No-op: mode configuration is not emulated
return ctx.SetReturn(0);
}
}
@@ -76,14 +76,6 @@ public static class MslFixedShaders
/// </summary>
public static string CreateDepthOnlyFragment() => MslTemplates.Render("depth_only_fragment");
/// <summary>
/// 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 <c>SpirvFixedShaders.CreateDetileCompute</c>. Entry point
/// "detile_cs"; buffers 0=tiled, 1=xTerm, 2=yTerm, 3=out, 4=DetileParams.
/// </summary>
public static string CreateDetileCompute() => MslTemplates.Render("detile_compute");
private static string Format(float value) =>
value.ToString("0.0######", CultureInfo.InvariantCulture) + "f";
}
@@ -1,70 +0,0 @@
#include <metal_stdlib>
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];
}
@@ -314,8 +314,7 @@ public static partial class Gen5SpirvTranslator
uint VectorType,
ImageComponentKind ComponentKind,
bool IsStorage,
bool Arrayed,
SpirvImageDim Dimension);
bool Arrayed);
private readonly record struct SpirvVertexInput(
uint Variable,
@@ -1002,15 +1001,11 @@ public static partial class Gen5SpirvTranslator
SpirvCapability.StorageImageExtendedFormats);
}
var dimension = binding.Control.Dimension == 2
? SpirvImageDim.Dim3D
: SpirvImageDim.Dim2D;
var isArrayed = dimension != SpirvImageDim.Dim3D &&
!isStorage &&
var isArrayed = !isStorage &&
Gen5ShaderTranslator.IsArrayedImageBinding(binding);
var imageType = _module.TypeImage(
componentType,
dimension,
SpirvImageDim.Dim2D,
depth: false,
arrayed: isArrayed,
multisampled: false,
@@ -1040,8 +1035,7 @@ public static partial class Gen5SpirvTranslator
_module.TypeVector(componentType, 4),
componentKind,
isStorage,
isArrayed,
dimension));
isArrayed));
_interfaces.Add(variable);
}
}
@@ -2265,39 +2259,20 @@ public static partial class Gen5SpirvTranslator
return false;
}
var memoryOpcode = control.UsesFlatAddress
? "Global" + instruction.Opcode["Flat".Length..]
: instruction.Opcode;
var vectorByteAddress = LoadV(control.VectorAddress);
if (control.UsesFlatAddress)
{
// FLAT instructions carry the complete 64-bit guest address
// in a VGPR pair. The scalar evaluator captures the buffer
// rooted at the inferred SGPR pair, so convert the low address
// dword back to a byte offset inside that binding. Subtraction
// modulo 2^32 also handles an address addition that carried
// into the high dword, while the bounded binding prevents an
// unrelated pointer from escaping into host storage.
vectorByteAddress = _module.AddInstruction(
SpirvOp.ISub,
_uintType,
vectorByteAddress,
LoadS(control.ScalarAddress));
}
var byteAddress = IAdd(
vectorByteAddress,
LoadV(control.VectorAddress),
UInt(unchecked((uint)control.OffsetBytes)));
byteAddress = ApplyGuestBufferByteBias(bindingIndex, byteAddress);
var dwordAddress = ShiftRightLogical(byteAddress, UInt(2));
if (memoryOpcode is "GlobalAtomicAdd" or "GlobalAtomicUMax")
if (instruction.Opcode is "GlobalAtomicAdd" or "GlobalAtomicUMax")
{
EmitExecConditional(() =>
{
EmitConditional(IsBufferWordInRange(bindingIndex, dwordAddress), () =>
{
var original = _module.AddInstruction(
memoryOpcode == "GlobalAtomicAdd"
instruction.Opcode == "GlobalAtomicAdd"
? SpirvOp.AtomicIAdd
: SpirvOp.AtomicUMax,
_uintType,
@@ -2314,12 +2289,12 @@ public static partial class Gen5SpirvTranslator
return true;
}
if (memoryOpcode.StartsWith("GlobalStore", StringComparison.Ordinal))
if (instruction.Opcode.StartsWith("GlobalStore", StringComparison.Ordinal))
{
EmitExecConditional(() =>
{
if (TryGetSubdwordStoreInfo(
memoryOpcode,
instruction.Opcode,
out var byteCount,
out var sourceShift))
{
@@ -2349,7 +2324,7 @@ public static partial class Gen5SpirvTranslator
}
if (TryGetSubdwordLoadInfo(
memoryOpcode,
instruction.Opcode,
out var loadByteCount,
out var signExtend,
out var d16,
@@ -3293,7 +3268,6 @@ public static partial class Gen5SpirvTranslator
var imageObject = Load(resource.ObjectType, resource.Variable);
if (instruction.Opcode == "ImageGetResinfo")
{
var sizeComponentCount = ImageCoordinateComponentCount(resource);
var queryImage = resource.IsStorage
? imageObject
: _module.AddInstruction(
@@ -3304,7 +3278,7 @@ public static partial class Gen5SpirvTranslator
resource.IsStorage
? SpirvOp.ImageQuerySize
: SpirvOp.ImageQuerySizeLod,
_module.TypeVector(_intType, sizeComponentCount),
_module.TypeVector(_intType, 2),
resource.IsStorage
? [queryImage]
: [queryImage, UInt(0)]);
@@ -3317,7 +3291,7 @@ public static partial class Gen5SpirvTranslator
}
uint value;
if (component < sizeComponentCount)
if (component < 2)
{
var signedValue = _module.AddInstruction(
SpirvOp.CompositeExtract,
@@ -3345,12 +3319,7 @@ public static partial class Gen5SpirvTranslator
return false;
}
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var coordinates = BuildIntegerCoordinates(
image,
0,
coordinateComponentCount);
var coordinates = BuildIntegerCoordinates(image, 0);
var components = new uint[4];
uint sourceIndex = 0;
for (var component = 0; component < components.Length; component++)
@@ -3386,14 +3355,13 @@ public static partial class Gen5SpirvTranslator
components);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, coordinateComponentCount),
_module.TypeVector(_intType, 2),
imageObject);
EmitBoundsCheckedImageWrite(
coordinates,
imageSize,
imageObject,
texel,
coordinateComponentCount);
texel);
return true;
}
@@ -3414,17 +3382,14 @@ public static partial class Gen5SpirvTranslator
}
var signed = resource.ComponentKind == ImageComponentKind.Sint;
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var atomicImageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, coordinateComponentCount),
_module.TypeVector(_intType, 2),
imageObject);
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
atomicImageSize,
coordinateComponentCount);
atomicImageSize);
EmitExecConditional(() =>
{
var pointer = _module.AddInstruction(
@@ -3468,17 +3433,14 @@ public static partial class Gen5SpirvTranslator
{
if (resource.IsStorage)
{
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, coordinateComponentCount),
_module.TypeVector(_intType, 2),
imageObject);
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
imageSize,
coordinateComponentCount);
imageSize);
sampled = _module.AddInstruction(
SpirvOp.ImageRead,
resource.VectorType,
@@ -3492,18 +3454,15 @@ public static partial class Gen5SpirvTranslator
SpirvOp.Image,
resource.ImageType,
imageObject);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySizeLod,
_module.TypeVector(_intType, coordinateComponentCount),
_module.TypeVector(_intType, 2),
fetchedImage,
UInt(mipLevel));
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
imageSize,
coordinateComponentCount);
imageSize);
sampled = _module.AddInstruction(
SpirvOp.ImageFetch,
resource.VectorType,
@@ -3535,19 +3494,12 @@ public static partial class Gen5SpirvTranslator
// lowering treated SAMPLE_D as body-first and consequently
// sampled gradients as coordinates in every captured
// derivative operation.
var spatialComponentCount =
ImageSpatialComponentCount(resource);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var addressCursor = 0;
var offset = 0u;
if (hasOffset)
{
addressCursor = AlignFullImageAddress(image, addressCursor);
offset = BuildImageOffset(
image,
addressCursor,
spatialComponentCount);
offset = BuildImageOffset(image, addressCursor);
addressCursor += ImageFullAddressSlots(image);
}
@@ -3571,33 +3523,26 @@ public static partial class Gen5SpirvTranslator
}
var gradientX = hasGradients
? BuildFloatCoordinates(
image,
addressCursor,
spatialComponentCount)
? BuildFloatCoordinates(image, addressCursor)
: 0u;
var gradientY = hasGradients
? BuildFloatCoordinates(
image,
addressCursor + (int)spatialComponentCount,
spatialComponentCount)
? BuildFloatCoordinates(image, addressCursor + 2)
: 0u;
if (hasGradients)
{
addressCursor += checked((int)(spatialComponentCount * 2));
addressCursor += 4;
}
var coordinates = BuildFloatCoordinates(
image,
addressCursor,
coordinateComponentCount);
var coordinates = resource.Arrayed
? BuildFloatArrayCoordinates(image, addressCursor)
: BuildFloatCoordinates(image, addressCursor);
var explicitLod = hasGradients || hasZeroLod || hasLod;
var lod = hasZeroLod
? Float(0)
: hasLod
? LoadImageFloatAddress(
image,
addressCursor + (int)coordinateComponentCount)
addressCursor + (resource.Arrayed ? 3 : 2))
: lodOrBias;
if (hasOffset)
{
@@ -3663,18 +3608,11 @@ public static partial class Gen5SpirvTranslator
instruction.Opcode.EndsWith("O", StringComparison.Ordinal);
var hasCompare =
instruction.Opcode.Contains("Gather4C", StringComparison.Ordinal);
var spatialComponentCount =
ImageSpatialComponentCount(resource);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var addressCursor = 0;
var offset = 0u;
if (hasOffset)
{
offset = BuildImageOffset(
image,
addressCursor,
spatialComponentCount);
offset = BuildImageOffset(image, addressCursor);
addressCursor += ImageFullAddressSlots(image);
}
@@ -3689,10 +3627,9 @@ public static partial class Gen5SpirvTranslator
addressCursor += ImageFullAddressSlots(image);
}
var coordinates = BuildFloatCoordinates(
image,
addressCursor,
coordinateComponentCount);
var coordinates = resource.Arrayed
? BuildFloatArrayCoordinates(image, addressCursor)
: BuildFloatCoordinates(image, addressCursor);
var operands = new List<uint>
{
imageObject,
@@ -3885,31 +3822,28 @@ public static partial class Gen5SpirvTranslator
});
}
private static uint ImageSpatialComponentCount(
SpirvImageResource resource) =>
resource.Dimension == SpirvImageDim.Dim3D ? 3u : 2u;
private static uint ImageCoordinateComponentCount(
SpirvImageResource resource) =>
resource.Arrayed ? 3u : ImageSpatialComponentCount(resource);
private uint BuildFloatCoordinates(
Gen5ImageControl image,
int start,
uint componentCount)
private uint BuildFloatCoordinates(Gen5ImageControl image, int start)
{
var components = new uint[checked((int)componentCount)];
for (var component = 0; component < components.Length; component++)
{
components[component] = LoadImageFloatAddress(
image,
start + component);
}
var x = LoadImageFloatAddress(image, start);
var y = LoadImageFloatAddress(image, start + 1);
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_module.TypeVector(_floatType, componentCount),
components);
_vec2Type,
x,
y);
}
private uint BuildFloatArrayCoordinates(Gen5ImageControl image, int start)
{
var x = LoadImageFloatAddress(image, start);
var y = LoadImageFloatAddress(image, start + 1);
var slice = LoadImageFloatAddress(image, start + 2);
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_vec3Type,
x,
y,
slice);
}
private static int ImageAddressRegister(
@@ -4014,49 +3948,47 @@ public static partial class Gen5SpirvTranslator
ShiftLeftLogical(BitwiseAnd(high, UInt(0xFFFF)), UInt(16)));
}
private uint BuildIntegerCoordinates(
Gen5ImageControl image,
int start,
uint componentCount)
private uint BuildIntegerCoordinates(Gen5ImageControl image, int start)
{
var components = new uint[checked((int)componentCount)];
for (var component = 0; component < components.Length; component++)
{
components[component] = Bitcast(
_intType,
LoadImageIntegerAddress(image, start + component));
}
var ivec2 = _module.TypeVector(_intType, 2);
var x = Bitcast(_intType, LoadImageIntegerAddress(image, start));
var y = Bitcast(_intType, LoadImageIntegerAddress(image, start + 1));
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_module.TypeVector(_intType, componentCount),
components);
ivec2,
x,
y);
}
private uint BuildClampedIntegerCoordinates(
Gen5ImageControl image,
int start,
uint imageSize,
uint componentCount)
uint imageSize)
{
var components = new uint[checked((int)componentCount)];
for (var component = 0; component < components.Length; component++)
{
components[component] = ClampSignedCoordinate(
Bitcast(
_intType,
LoadImageIntegerAddress(image, start + component)),
_module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
(uint)component));
}
var ivec2 = _module.TypeVector(_intType, 2);
var x = ClampSignedCoordinate(
Bitcast(
_intType,
LoadImageIntegerAddress(image, start)),
_module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
0));
var y = ClampSignedCoordinate(
Bitcast(
_intType,
LoadImageIntegerAddress(image, start + 1)),
_module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
1));
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_module.TypeVector(_intType, componentCount),
components);
ivec2,
x,
y);
}
private uint ClampSignedCoordinate(uint value, uint extent)
@@ -4095,47 +4027,69 @@ public static partial class Gen5SpirvTranslator
uint coordinates,
uint imageSize,
uint imageObject,
uint texel,
uint coordinateComponentCount)
uint texel)
{
var x = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
coordinates,
0);
var y = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
coordinates,
1);
var width = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
0);
var height = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
1);
var zero = _module.Constant(_intType, 0);
var inRange = Load(_boolType, _exec);
for (uint component = 0;
component < coordinateComponentCount;
component++)
{
var coordinate = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
coordinates,
component);
var extent = _module.AddInstruction(
SpirvOp.CompositeExtract,
_intType,
imageSize,
component);
var nonNegative = _module.AddInstruction(
SpirvOp.SGreaterThanEqual,
_boolType,
coordinate,
zero);
var belowExtent = _module.AddInstruction(
SpirvOp.SLessThan,
_boolType,
coordinate,
extent);
var componentInRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
nonNegative,
belowExtent);
inRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
inRange,
componentInRange);
}
var xNonNegative = _module.AddInstruction(
SpirvOp.SGreaterThanEqual,
_boolType,
x,
zero);
var yNonNegative = _module.AddInstruction(
SpirvOp.SGreaterThanEqual,
_boolType,
y,
zero);
var xInRange = _module.AddInstruction(
SpirvOp.SLessThan,
_boolType,
x,
width);
var yInRange = _module.AddInstruction(
SpirvOp.SLessThan,
_boolType,
y,
height);
var lowerInRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
xNonNegative,
yNonNegative);
var upperInRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
xInRange,
yInRange);
var inRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
lowerInRange,
upperInRange);
inRange = _module.AddInstruction(
SpirvOp.LogicalAnd,
_boolType,
Load(_boolType, _exec),
inRange);
var writeLabel = _module.AllocateId();
var mergeLabel = _module.AllocateId();
_module.AddStatement(SpirvOp.SelectionMerge, mergeLabel, 0);
@@ -4154,30 +4108,30 @@ public static partial class Gen5SpirvTranslator
_module.AddLabel(mergeLabel);
}
private uint BuildImageOffset(
Gen5ImageControl image,
int component,
uint componentCount)
private uint BuildImageOffset(Gen5ImageControl image, int component)
{
var ivec2 = _module.TypeVector(_intType, 2);
var packed = Bitcast(
_intType,
LoadV(image.GetAddressRegister(
ImageAddressRegister(image, component))));
var components = new uint[checked((int)componentCount)];
for (var index = 0; index < components.Length; index++)
{
components[index] = _module.AddInstruction(
SpirvOp.BitFieldSExtract,
_intType,
packed,
UInt((uint)(index * 8)),
UInt(6));
}
var x = _module.AddInstruction(
SpirvOp.BitFieldSExtract,
_intType,
packed,
UInt(0),
UInt(6));
var y = _module.AddInstruction(
SpirvOp.BitFieldSExtract,
_intType,
packed,
UInt(8),
UInt(6));
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_module.TypeVector(_intType, componentCount),
components);
ivec2,
x,
y);
}
private uint ApplyDynamicSampleOffset(
@@ -4187,16 +4141,7 @@ public static partial class Gen5SpirvTranslator
uint texelOffset,
uint lod)
{
var spatialComponentCount = ImageSpatialComponentCount(resource);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var spatialIntegerType =
_module.TypeVector(_intType, spatialComponentCount);
var spatialFloatType =
_module.TypeVector(_floatType, spatialComponentCount);
var queryComponentCount = resource.Arrayed
? coordinateComponentCount
: spatialComponentCount;
var ivec2 = _module.TypeVector(_intType, 2);
var image = _module.AddInstruction(
SpirvOp.Image,
resource.ImageType,
@@ -4218,14 +4163,14 @@ public static partial class Gen5SpirvTranslator
signedLod);
var size = _module.AddInstruction(
SpirvOp.ImageQuerySizeLod,
_module.TypeVector(_intType, queryComponentCount),
resource.Arrayed ? _module.TypeVector(_intType, 3) : ivec2,
image,
clampedLod);
if (resource.Arrayed)
{
size = _module.AddInstruction(
SpirvOp.VectorShuffle,
spatialIntegerType,
ivec2,
size,
size,
0u,
@@ -4234,27 +4179,27 @@ public static partial class Gen5SpirvTranslator
var sizeFloat = _module.AddInstruction(
SpirvOp.ConvertSToF,
spatialFloatType,
_vec2Type,
size);
var offsetFloat = _module.AddInstruction(
SpirvOp.ConvertSToF,
spatialFloatType,
_vec2Type,
texelOffset);
var normalizedOffset = _module.AddInstruction(
SpirvOp.FDiv,
spatialFloatType,
_vec2Type,
offsetFloat,
sizeFloat);
if (!resource.Arrayed)
{
return _module.AddInstruction(
SpirvOp.FAdd,
spatialFloatType,
_vec2Type,
coordinates,
normalizedOffset);
}
var arrayOffset = _module.AddInstruction(
var offsetVec3 = _module.AddInstruction(
SpirvOp.CompositeConstruct,
_vec3Type,
_module.AddInstruction(
@@ -4272,7 +4217,7 @@ public static partial class Gen5SpirvTranslator
SpirvOp.FAdd,
_vec3Type,
coordinates,
arrayOffset);
offsetVec3);
}
private bool TryEmitExport(
@@ -260,233 +260,4 @@ public static class SpirvFixedShaders
module.AddExecutionMode(main, SpirvExecutionMode.OriginUpperLeft);
return module.Build();
}
/// <summary>
/// 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 <c>GnmTiling.GetDetileParams</c> so it is bit-identical
/// to the CPU fallback for both supported equation families:
/// <code>
/// z = layer;
/// inBlock = equation == BlockTable // modes 1/4/8
/// ? blockTable[(y % blockHeight) * blockWidth + (x % blockWidth)]
/// : xTerm[x &amp; xMask] ^ yTerm[y &amp; 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];
/// </code>
/// Each array slice is an independently tiled 2D surface; the caller packs the
/// slices contiguously in the tiled buffer (stride <c>srcSliceElements</c>) and
/// the output ends up layer-major, matching a single multi-layer
/// buffer-&gt;image copy. For a non-arrayed texture the caller dispatches a
/// single Z layer with <c>srcSliceElements</c> 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.
/// </summary>
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();
}
}
+1 -2
View File
@@ -192,8 +192,7 @@ public sealed record Gen5GlobalMemoryControl(
uint ScalarAddress,
int OffsetBytes,
bool Glc,
bool Slc,
bool UsesFlatAddress = false) : Gen5InstructionControl;
bool Slc) : Gen5InstructionControl;
public sealed record Gen5BufferMemoryControl(
uint DwordCount,
@@ -4,7 +4,6 @@
using SharpEmu.HLE;
using System.Buffers;
using System.Buffers.Binary;
using System.Collections.Concurrent;
using System.Diagnostics;
using System.Numerics;
@@ -36,14 +35,6 @@ public static class Gen5ShaderScalarEvaluator
StringComparison.Ordinal);
private static readonly object _scalarFallbackTraceGate = new();
private static readonly HashSet<(ulong Shader, uint Pc)> _tracedScalarFallbacks = [];
// Shaders whose empty SRT/EUD caused a null-base scalar pointer load.
// Host submit of those translations has lost the Vulkan device; Agc skips
// them before QueueSubmit.
private static readonly ConcurrentDictionary<ulong, byte> _emptySrtScalarPointerFallbacks =
new();
public static bool WasEmptySrtScalarPointerFallback(ulong shaderAddress) =>
_emptySrtScalarPointerFallbacks.ContainsKey(shaderAddress);
// Uniform forward branches select material/resource bodies that remain
// statically present in the translated shader. Discover the skipped body's
@@ -358,10 +349,7 @@ public static class Gen5ShaderScalarEvaluator
if (globalMemory.ScalarAddress >= ScalarRegisterCount - 1)
{
error =
$"{(globalMemory.UsesFlatAddress
? "flat-address-base-unresolved"
: "global-address-register-range")} " +
$"pc=0x{instruction.Pc:X} " +
$"global-address-register-range pc=0x{instruction.Pc:X} " +
$"s{globalMemory.ScalarAddress}";
return false;
}
@@ -376,18 +364,11 @@ public static class Gen5ShaderScalarEvaluator
}
var key = (globalMemory.ScalarAddress, baseAddress);
var writable =
instruction.Opcode.StartsWith(
var writable = instruction.Opcode.StartsWith(
"GlobalStore",
StringComparison.Ordinal) ||
instruction.Opcode.StartsWith(
"GlobalAtomic",
StringComparison.Ordinal) ||
instruction.Opcode.StartsWith(
"FlatStore",
StringComparison.Ordinal) ||
instruction.Opcode.StartsWith(
"FlatAtomic",
StringComparison.Ordinal);
if (globalMemoryByAddress.TryGetValue(key, out var existingBinding))
{
@@ -2120,15 +2101,6 @@ public static class Gen5ShaderScalarEvaluator
$"dynamic={dynamicOffset} definitions=[{string.Join(';', definitions)}] " +
$"user_data=[{userData}] metadata=" +
$"{(state.Metadata is null ? "missing" : $"srt={state.Metadata.ShaderResourceTableSizeDwords},eud={state.Metadata.ExtendedUserDataSizeDwords}")}");
if (baseAddress == 0 &&
state.Metadata is
{
ShaderResourceTableSizeDwords: 0,
ExtendedUserDataSizeDwords: 0,
})
{
_emptySrtScalarPointerFallbacks.TryAdd(state.Program.Address, 0);
}
}
[Conditional("DEBUG")]
@@ -305,23 +305,6 @@ public static class Gen5ShaderTranslator
count |= 0x20;
}
// Primary SH defaults leave SPI_SHADER_PGM_RSRC2_PS at 0. Draws that
// still wrote USER_DATA_n via SetShReg would otherwise translate with
// an empty SRT window (Astro title PS → Address-0 descriptors →
// device lost). Recover the window from contiguous live registers.
if (count == 0 &&
userDataBaseRegister is not ComputeUserDataRegister)
{
var probed = 0;
while (probed < MaximumHardwareUserSgprs &&
shaderRegisters.ContainsKey(userDataBaseRegister + (uint)probed))
{
probed++;
}
count = probed;
}
if (userDataBaseRegister is not (PsUserDataRegister or
VsUserDataRegister or
GsUserDataRegister or
@@ -471,15 +454,7 @@ public static class Gen5ShaderTranslator
}
}
var instruction = CreateInstruction(pc, encoding, name, words);
if (instruction.Control is Gen5GlobalMemoryControl
{
UsesFlatAddress: true,
})
{
instruction = ResolveFlatAddressBase(instructions, instruction);
}
instructions.Add(instruction);
instructions.Add(CreateInstruction(pc, encoding, name, words));
instructionCount++;
pc += sizeDwords * sizeof(uint);
@@ -1403,42 +1378,35 @@ public static class Gen5ShaderTranslator
var opcode = (word >> 18) & 0x7F;
sizeDwords = 2;
error = string.Empty;
var prefix = segment switch
{
0x0 => "Flat",
0x2 => "Global",
_ => string.Empty,
};
var suffix = opcode switch
{
0x08 => "LoadUbyte",
0x09 => "LoadSbyte",
0x0A => "LoadUshort",
0x0B => "LoadSshort",
0x0C => "LoadDword",
0x0D => "LoadDwordx2",
0x0E => "LoadDwordx4",
0x0F => "LoadDwordx3",
0x18 => "StoreByte",
0x19 => "StoreByteD16Hi",
0x1A => "StoreShort",
0x1B => "StoreShortD16Hi",
0x1C => "StoreDword",
0x1D => "StoreDwordx2",
0x1E => "StoreDwordx4",
0x1F => "StoreDwordx3",
0x20 => "LoadUbyteD16",
0x21 => "LoadUbyteD16Hi",
0x22 => "LoadSbyteD16",
0x23 => "LoadSbyteD16Hi",
0x24 => "LoadShortD16",
0x25 => "LoadShortD16Hi",
0x32 => "AtomicAdd",
0x38 => "AtomicUMax",
_ => string.Empty,
};
name = prefix.Length != 0 && suffix.Length != 0
? prefix + suffix
name = segment == 0x2
? opcode switch
{
0x08 => "GlobalLoadUbyte",
0x09 => "GlobalLoadSbyte",
0x0A => "GlobalLoadUshort",
0x0B => "GlobalLoadSshort",
0x0C => "GlobalLoadDword",
0x0D => "GlobalLoadDwordx2",
0x0E => "GlobalLoadDwordx4",
0x0F => "GlobalLoadDwordx3",
0x18 => "GlobalStoreByte",
0x19 => "GlobalStoreByteD16Hi",
0x1A => "GlobalStoreShort",
0x1B => "GlobalStoreShortD16Hi",
0x1C => "GlobalStoreDword",
0x1D => "GlobalStoreDwordx2",
0x1E => "GlobalStoreDwordx4",
0x1F => "GlobalStoreDwordx3",
0x20 => "GlobalLoadUbyteD16",
0x21 => "GlobalLoadUbyteD16Hi",
0x22 => "GlobalLoadSbyteD16",
0x23 => "GlobalLoadSbyteD16Hi",
0x24 => "GlobalLoadShortD16",
0x25 => "GlobalLoadShortD16Hi",
0x32 => "GlobalAtomicAdd",
0x38 => "GlobalAtomicUMax",
_ => string.Empty,
}
: string.Empty;
return FinishDecode(
@@ -1661,82 +1629,6 @@ public static class Gen5ShaderTranslator
_ => false,
};
private static Gen5ShaderInstruction ResolveFlatAddressBase(
IReadOnlyList<Gen5ShaderInstruction> precedingInstructions,
Gen5ShaderInstruction instruction)
{
if (instruction.Control is not Gen5GlobalMemoryControl
{
UsesFlatAddress: true,
} control ||
!TryFindVectorDefinition(
precedingInstructions,
control.VectorAddress,
out var lowDefinition) ||
!TryFindVectorDefinition(
precedingInstructions,
control.VectorAddress + 1,
out var highDefinition))
{
return instruction;
}
foreach (var lowSource in lowDefinition.Sources)
{
if (lowSource.Kind != Gen5OperandKind.ScalarRegister)
{
continue;
}
foreach (var highSource in highDefinition.Sources)
{
if (highSource.Kind != Gen5OperandKind.ScalarRegister ||
highSource.Value != lowSource.Value + 1)
{
continue;
}
return instruction with
{
Sources =
[
.. instruction.Sources,
Gen5Operand.Scalar(lowSource.Value),
],
Control = control with
{
ScalarAddress = lowSource.Value,
},
};
}
}
return instruction;
}
private static bool TryFindVectorDefinition(
IReadOnlyList<Gen5ShaderInstruction> instructions,
uint register,
out Gen5ShaderInstruction definition)
{
for (var index = instructions.Count - 1; index >= 0; index--)
{
var candidate = instructions[index];
foreach (var destination in candidate.Destinations)
{
if (destination.Kind == Gen5OperandKind.VectorRegister &&
destination.Value == register)
{
definition = candidate;
return true;
}
}
}
definition = default!;
return false;
}
private static Gen5ShaderInstruction CreateInstruction(
uint pc,
Gen5ShaderEncoding encoding,
@@ -2148,13 +2040,7 @@ public static class Gen5ShaderTranslator
var vectorAddress = extra & 0xFF;
var vectorData = (extra >> 8) & 0xFF;
var scalarAddress = (extra >> 16) & 0x7F;
var usesFlatAddress = opcode.StartsWith(
"Flat",
StringComparison.Ordinal);
var memoryOpcode = usesFlatAddress
? "Global" + opcode["Flat".Length..]
: opcode;
var dwordCount = memoryOpcode switch
var dwordCount = opcode switch
{
"GlobalLoadUbyte" or
"GlobalLoadSbyte" or
@@ -2182,20 +2068,12 @@ public static class Gen5ShaderTranslator
"GlobalStoreDwordx4" => 4u,
_ => 0u,
};
sources = usesFlatAddress
?
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Vector(vectorAddress + 1),
]
:
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Scalar(scalarAddress),
];
destinations = memoryOpcode.StartsWith(
"GlobalLoad",
StringComparison.Ordinal)
sources =
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Scalar(scalarAddress),
];
destinations = opcode.StartsWith("GlobalLoad", StringComparison.Ordinal)
? Enumerable
.Range((int)vectorData, checked((int)dwordCount))
.Select(index => Gen5Operand.Vector((uint)index))
@@ -2205,11 +2083,10 @@ public static class Gen5ShaderTranslator
dwordCount,
vectorAddress,
vectorData,
usesFlatAddress ? uint.MaxValue : scalarAddress,
scalarAddress,
SignExtend(word & 0x1FFF, 13),
((word >> 16) & 1) != 0,
((word >> 17) & 1) != 0,
usesFlatAddress);
((word >> 17) & 1) != 0);
break;
}
case Gen5ShaderEncoding.Mubuf:
@@ -1,106 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using SharpEmu.Libs.Agc;
using SharpEmu.Libs.Gpu;
using Xunit;
namespace SharpEmu.Libs.Tests.Agc;
public sealed class AgcTextureTransportTests
{
[Theory]
[InlineData(10u, 4u, 4u)]
[InlineData(10u, 0u, 1u)]
[InlineData(9u, 4u, 1u)]
[InlineData(13u, 4u, 1u)]
public void GetTextureVolumeDepth_OnlyUsesDescriptorDepthFor3D(
uint type,
uint descriptorDepth,
uint expectedDepth)
{
Assert.Equal(
expectedDepth,
AgcExports.GetTextureVolumeDepth(type, descriptorDepth));
}
[Fact]
public void GetTextureByteCount_MultipliesUncompressedVolumeDepth()
{
Assert.Equal(
4UL * 8 * 6 * 5,
AgcExports.GetTextureByteCount(
format: 10,
width: 8,
height: 6,
depth: 5));
}
[Fact]
public void GetTextureByteCount_MultipliesBlockCompressedVolumeDepth()
{
// Format 169 uses one eight-byte BC block for each 4x4 texel block.
Assert.Equal(
2UL * 2 * 8 * 3,
AgcExports.GetTextureByteCount(
format: 169,
width: 7,
height: 5,
depth: 3));
}
[Fact]
public void GetTextureByteCount_LeavesTwoDimensionalSizingUnchanged()
{
Assert.Equal(
AgcExports.GetTextureByteCount(10, 8, 6),
AgcExports.GetTextureByteCount(10, 8, 6, depth: 1));
Assert.Equal(
AgcExports.GetTextureByteCount(10, 8, 6),
AgcExports.GetTextureByteCount(10, 8, 6, depth: 0));
}
[Fact]
public void GuestDrawTexture_CarriesRawTypeAndNormalizedDepth()
{
var texture = new GuestDrawTexture(
Address: 0x1234,
Width: 8,
Height: 6,
Format: 10,
NumberType: 0,
RgbaPixels: [],
IsFallback: false,
IsStorage: false,
Type: 10,
Depth: 5);
Assert.Equal(10u, texture.Type);
Assert.Equal(5u, texture.Depth);
}
[Fact]
public void TextureContentIdentity_DistinguishesTypeAndDepth()
{
var twoDimensional = CreateIdentity(type: 9, depth: 1);
var threeDimensional = CreateIdentity(type: 10, depth: 1);
var deeperThreeDimensional = CreateIdentity(type: 10, depth: 5);
Assert.NotEqual(twoDimensional, threeDimensional);
Assert.NotEqual(threeDimensional, deeperThreeDimensional);
}
private static TextureContentIdentity CreateIdentity(uint type, uint depth) =>
new(
Address: 0x1234,
Width: 8,
Height: 6,
Format: 10,
NumberType: 0,
DstSelect: 0xFAC,
TileMode: 0,
Pitch: 8,
Sampler: default,
Type: type,
Depth: depth);
}
@@ -1,102 +0,0 @@
// 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);
}
}
@@ -83,152 +83,4 @@ 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;
}
}
@@ -1,231 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Buffers.Binary;
using SharpEmu.HLE;
using SharpEmu.HLE.Host;
using SharpEmu.Libs.Audio;
using Xunit;
namespace SharpEmu.Libs.Tests.Audio;
[CollectionDefinition("AudioOutState", DisableParallelization = true)]
public sealed class AudioOutStateCollection
{
public const string Name = "AudioOutState";
}
[Collection(AudioOutStateCollection.Name)]
public sealed class AudioOutExportsTests : IDisposable
{
private const ulong MemoryBase = 0x1_0000_0000;
private const int MemorySize = 0x4000;
private const ulong ParameterAddress = MemoryBase + 0x100;
private const ulong FirstSourceAddress = MemoryBase + 0x1000;
private const ulong SecondSourceAddress = MemoryBase + 0x2000;
private readonly FakeCpuMemory _memory = new(MemoryBase, MemorySize);
private readonly CpuContext _ctx;
private readonly List<RecordingAudioStream> _streams = [];
public AudioOutExportsTests()
{
AudioOutExports.ResetForTests();
AudioOutExports.SetStreamFactoryForTests(_ =>
{
var stream = new RecordingAudioStream();
_streams.Add(stream);
return stream;
});
_ctx = new CpuContext(_memory, Generation.Gen5);
}
[Fact]
public void Outputs_StagesAndSubmitsSingleStereoPort()
{
var handle = OpenPort(bufferLength: 2);
Span<byte> source = stackalloc byte[8];
BinaryPrimitives.WriteInt16LittleEndian(source, -32768);
BinaryPrimitives.WriteInt16LittleEndian(source[2..], 32767);
BinaryPrimitives.WriteInt16LittleEndian(source[4..], -1234);
BinaryPrimitives.WriteInt16LittleEndian(source[6..], 5678);
Assert.True(_memory.TryWrite(FirstSourceAddress, source));
WriteDescriptor(0, handle, FirstSourceAddress);
var result = Submit(outputCount: 1);
Assert.Equal(2, result);
Assert.Equal(2UL, _ctx[CpuRegister.Rax]);
Assert.Equal(source.ToArray(), Assert.Single(_streams[0].Submissions));
}
[Fact]
public void Outputs_SubmitsEveryPortInTheBatch()
{
var firstHandle = OpenPort(bufferLength: 2);
var secondHandle = OpenPort(bufferLength: 2);
byte[] firstSource = [1, 0, 2, 0, 3, 0, 4, 0];
byte[] secondSource = [5, 0, 6, 0, 7, 0, 8, 0];
Assert.True(_memory.TryWrite(FirstSourceAddress, firstSource));
Assert.True(_memory.TryWrite(SecondSourceAddress, secondSource));
// Reverse handle order to exercise canonical lock ordering without
// changing which guest buffer belongs to each port.
WriteDescriptor(0, secondHandle, FirstSourceAddress);
WriteDescriptor(1, firstHandle, SecondSourceAddress);
var result = Submit(outputCount: 2);
Assert.Equal(2, result);
Assert.Equal(secondSource, Assert.Single(_streams[0].Submissions));
Assert.Equal(firstSource, Assert.Single(_streams[1].Submissions));
}
[Fact]
public void Outputs_AcceptsNullBufferAsSynchronizationOnly()
{
var handle = OpenPort(bufferLength: 2);
WriteDescriptor(0, handle, sourceAddress: 0);
var result = Submit(outputCount: 1);
Assert.Equal(2, result);
Assert.Empty(_streams[0].Submissions);
}
[Fact]
public void Outputs_FaultInLaterBufferDoesNotPartiallySubmit()
{
var firstHandle = OpenPort(bufferLength: 2);
var secondHandle = OpenPort(bufferLength: 2);
Assert.True(_memory.TryWrite(FirstSourceAddress, new byte[8]));
WriteDescriptor(0, firstHandle, FirstSourceAddress);
WriteDescriptor(1, secondHandle, MemoryBase + MemorySize);
var result = Submit(outputCount: 2);
Assert.Equal(AudioOutExports.AudioOutErrorInvalidPointer, result);
Assert.All(_streams, stream => Assert.Empty(stream.Submissions));
}
[Fact]
public void Outputs_RejectsDuplicateHandlesBeforeSubmission()
{
var handle = OpenPort(bufferLength: 2);
Assert.True(_memory.TryWrite(FirstSourceAddress, new byte[8]));
Assert.True(_memory.TryWrite(SecondSourceAddress, new byte[8]));
WriteDescriptor(0, handle, FirstSourceAddress);
WriteDescriptor(1, handle, SecondSourceAddress);
var result = Submit(outputCount: 2);
Assert.Equal(AudioOutExports.AudioOutErrorInvalidPort, result);
Assert.Empty(_streams[0].Submissions);
}
[Fact]
public void Outputs_RejectsPortsWithDifferentBufferLengths()
{
var firstHandle = OpenPort(bufferLength: 2);
var secondHandle = OpenPort(bufferLength: 4);
WriteDescriptor(0, firstHandle, FirstSourceAddress);
WriteDescriptor(1, secondHandle, SecondSourceAddress);
var result = Submit(outputCount: 2);
Assert.Equal(AudioOutExports.AudioOutErrorInvalidSize, result);
Assert.All(_streams, stream => Assert.Empty(stream.Submissions));
}
[Fact]
public void Outputs_RejectsUnknownHandle()
{
WriteDescriptor(0, 99, FirstSourceAddress);
Assert.Equal(
AudioOutExports.AudioOutErrorInvalidPort,
Submit(outputCount: 1));
}
[Theory]
[InlineData(0u)]
[InlineData(26u)]
public void Outputs_RejectsInvalidOutputCount(uint outputCount)
{
Assert.Equal(
AudioOutExports.AudioOutErrorPortFull,
Submit(outputCount));
}
[Fact]
public void Outputs_RejectsNullOrUnreadableParameterArray()
{
Assert.Equal(
AudioOutExports.AudioOutErrorInvalidPointer,
Submit(outputCount: 1, parameterAddress: 0));
Assert.Equal(
AudioOutExports.AudioOutErrorInvalidPointer,
Submit(outputCount: 1, parameterAddress: MemoryBase + MemorySize));
}
[Fact]
public void OutputsExportRegistersForBothGenerations()
{
foreach (var generation in new[] { Generation.Gen4, Generation.Gen5 })
{
var manager = new ModuleManager();
manager.RegisterExports(
SharpEmu.Generated.SysAbiExportRegistry.CreateExports(generation));
Assert.True(manager.TryGetExport("w3PdaSTSwGE", out var export));
Assert.Equal("sceAudioOutOutputs", export.Name);
Assert.Equal("libSceAudioOut", export.LibraryName);
}
}
public void Dispose() => AudioOutExports.ResetForTests();
private int OpenPort(uint bufferLength)
{
_ctx[CpuRegister.Rdi] = 1;
_ctx[CpuRegister.Rsi] = 0;
_ctx[CpuRegister.Rdx] = 0;
_ctx[CpuRegister.Rcx] = bufferLength;
_ctx[CpuRegister.R8] = 48000;
_ctx[CpuRegister.R9] = 1;
return AudioOutExports.AudioOutOpen(_ctx);
}
private int Submit(uint outputCount, ulong parameterAddress = ParameterAddress)
{
_ctx[CpuRegister.Rdi] = parameterAddress;
_ctx[CpuRegister.Rsi] = outputCount;
return AudioOutExports.AudioOutOutputs(_ctx);
}
private void WriteDescriptor(int index, int handle, ulong sourceAddress)
{
Span<byte> descriptor = stackalloc byte[16];
descriptor.Clear();
BinaryPrimitives.WriteInt32LittleEndian(descriptor, handle);
BinaryPrimitives.WriteUInt64LittleEndian(descriptor[8..], sourceAddress);
Assert.True(_memory.TryWrite(
ParameterAddress + unchecked((ulong)(index * descriptor.Length)),
descriptor));
}
private sealed class RecordingAudioStream : IHostAudioStream
{
public List<byte[]> Submissions { get; } = [];
public bool Submit(ReadOnlySpan<byte> stereoPcm16)
{
Submissions.Add(stereoPcm16.ToArray());
return true;
}
public void Dispose()
{
}
}
}
@@ -83,9 +83,7 @@ public sealed class BinkFramePlaybackTests
public uint Height => 1;
// Keep frame boundaries far enough apart that a loaded CI runner cannot
// skip an expected frame between polling iterations.
public uint FramesPerSecondNumerator => 2;
public uint FramesPerSecondNumerator => 20;
public uint FramesPerSecondDenominator => 1;
@@ -11,11 +11,6 @@ public sealed class SystemServiceExportsTests
{
private const ulong MemoryBase = 0x1_0000_0000;
public SystemServiceExportsTests()
{
SystemServiceExports.ResetForTests();
}
[Fact]
public void GetNoticeScreenSkipFlagWritesOneByteAtMemoryBoundary()
{
@@ -31,23 +26,4 @@ public sealed class SystemServiceExportsTests
Assert.True(memory.TryRead(MemoryBase, flag));
Assert.Equal(0, flag[0]);
}
[Fact]
public void SetNoticeScreenSkipFlagRoundTripsThroughGetter()
{
var memory = new FakeCpuMemory(MemoryBase, 2);
var context = new CpuContext(memory, Generation.Gen5)
{
[CpuRegister.Rdi] = 1,
};
Assert.Equal(0, SystemServiceExports.SystemServiceSetNoticeScreenSkipFlag(context));
context[CpuRegister.Rdi] = MemoryBase;
Assert.Equal(0, SystemServiceExports.SystemServiceGetNoticeScreenSkipFlag(context));
Span<byte> flag = stackalloc byte[1];
Assert.True(memory.TryRead(MemoryBase, flag));
Assert.Equal(1, flag[0]);
}
}
@@ -36,23 +36,4 @@ public sealed class VulkanGuestImageByteCountTests
expected,
VulkanVideoPresenter.GetGuestImageByteCount(format, width, height));
}
[Theory]
[InlineData(10u, 8u, 4u, 3u, 384UL)]
[InlineData(169u, 5u, 5u, 7u, 224UL)]
public void MultipliesSurfaceSizeByVolumeDepth(
uint format,
uint width,
uint height,
uint depth,
ulong expected)
{
Assert.Equal(
expected,
VulkanVideoPresenter.GetGuestImageByteCount(
format,
width,
height,
depth));
}
}
@@ -1,49 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using Silk.NET.Vulkan;
using SharpEmu.Libs.VideoOut;
using Xunit;
namespace SharpEmu.Libs.Tests.VideoOut;
public sealed class VulkanGuestImageTypeTests
{
[Fact]
public void ThreeDimensionalDescriptorsMapToVolumeImageAndViewTypes()
{
Assert.Equal(
ImageType.Type3D,
VulkanVideoPresenter.GetGuestTextureImageType(
VulkanVideoPresenter.Gen5TextureType3D));
Assert.Equal(
ImageViewType.Type3D,
VulkanVideoPresenter.GetGuestTextureViewType(
VulkanVideoPresenter.Gen5TextureType3D,
arrayedView: true));
Assert.Equal(
7u,
VulkanVideoPresenter.GetGuestTextureDepth(
VulkanVideoPresenter.Gen5TextureType3D,
7));
}
[Fact]
public void TwoDimensionalArraysKeepLayersSeparateFromImageDepth()
{
Assert.Equal(
ImageType.Type2D,
VulkanVideoPresenter.GetGuestTextureImageType(
VulkanVideoPresenter.Gen5TextureType2D));
Assert.Equal(
ImageViewType.Type2DArray,
VulkanVideoPresenter.GetGuestTextureViewType(
VulkanVideoPresenter.Gen5TextureType2D,
arrayedView: true));
Assert.Equal(
1u,
VulkanVideoPresenter.GetGuestTextureDepth(
VulkanVideoPresenter.Gen5TextureType2D,
7));
}
}
@@ -1,179 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Buffers.Binary;
using SharpEmu.HLE;
using SharpEmu.ShaderCompiler;
using SharpEmu.ShaderCompiler.Vulkan;
using Xunit;
namespace SharpEmu.ShaderCompiler.Tests;
public sealed class Gen5FlatMemoryTests
{
private const ulong ShaderAddress = 0x1_0000_0000;
private const uint SEndpgm = 0xBF810000;
[Fact]
public void FlatLoadUbyteInfersScalarBaseAndCompiles()
{
var memory = new TestCpuMemory(ShaderAddress, 0x4000);
uint[] words =
[
// v_add_co_u32 v1, vcc_lo, s12, v6
0xD70F6A01,
0x00020C0C,
// v_add_co_ci_u32_sdwa v2, vcc_lo, 0, s13, vcc_lo
0x50041AF9,
0x86860680,
// flat_load_ubyte v0, v[1:2]
0xDC200000,
0x007D0001,
SEndpgm,
];
var shader = new byte[words.Length * sizeof(uint)];
for (var index = 0; index < words.Length; index++)
{
BinaryPrimitives.WriteUInt32LittleEndian(
shader.AsSpan(index * sizeof(uint)),
words[index]);
}
Assert.True(memory.TryWrite(ShaderAddress, shader));
var ctx = new CpuContext(memory, Generation.Gen5);
Assert.True(
Gen5ShaderTranslator.TryDecodeProgram(
ctx,
ShaderAddress,
out var program,
out var decodeError),
decodeError);
var instruction = Assert.Single(
program.Instructions,
item => item.Opcode == "FlatLoadUbyte");
var control = Assert.IsType<Gen5GlobalMemoryControl>(
instruction.Control);
Assert.True(control.UsesFlatAddress);
Assert.Equal(1u, control.VectorAddress);
Assert.Equal(0u, control.VectorData);
Assert.Equal(12u, control.ScalarAddress);
Assert.Equal(
[
Gen5Operand.Vector(1),
Gen5Operand.Vector(2),
Gen5Operand.Scalar(12),
],
instruction.Sources);
uint[] userData =
[
unchecked((uint)ShaderAddress),
unchecked((uint)(ShaderAddress >> 32)),
];
var state = new Gen5ShaderState(
program,
userData,
null,
UserDataScalarRegisterBase: 12);
Assert.True(
Gen5ShaderScalarEvaluator.TryEvaluate(
ctx,
state,
out var evaluation,
out var evaluationError),
evaluationError);
var binding = Assert.Single(evaluation.GlobalMemoryBindings);
Assert.Equal(12u, binding.ScalarAddress);
Assert.Contains(instruction.Pc, binding.InstructionPcs);
Assert.True(
Gen5SpirvTranslator.TryCompileComputeShader(
state,
evaluation,
1,
1,
1,
out var compiled,
out var compileError),
compileError);
Assert.Contains(
(ushort)SpirvOp.ISub,
ReadSpirvOpcodes(compiled.Spirv));
}
private static IReadOnlyList<ushort> ReadSpirvOpcodes(byte[] spirv)
{
Assert.Equal(0, spirv.Length % sizeof(uint));
Assert.True(spirv.Length >= 5 * sizeof(uint));
Assert.Equal(
0x07230203u,
BinaryPrimitives.ReadUInt32LittleEndian(spirv));
var opcodes = new List<ushort>();
for (var offset = 5 * sizeof(uint); offset < spirv.Length;)
{
var instruction =
BinaryPrimitives.ReadUInt32LittleEndian(spirv.AsSpan(offset));
var wordCount = checked((int)(instruction >> 16));
Assert.InRange(
wordCount,
1,
(spirv.Length - offset) / sizeof(uint));
opcodes.Add((ushort)instruction);
offset += wordCount * sizeof(uint);
}
return opcodes;
}
private sealed class TestCpuMemory(ulong baseAddress, int size) : ICpuMemory
{
private readonly byte[] _storage = new byte[size];
public bool TryRead(ulong virtualAddress, Span<byte> destination)
{
if (!TryResolve(virtualAddress, destination.Length, out var offset))
{
return false;
}
_storage.AsSpan(offset, destination.Length).CopyTo(destination);
return true;
}
public bool TryWrite(
ulong virtualAddress,
ReadOnlySpan<byte> source)
{
if (!TryResolve(virtualAddress, source.Length, out var offset))
{
return false;
}
source.CopyTo(_storage.AsSpan(offset, source.Length));
return true;
}
private bool TryResolve(
ulong virtualAddress,
int length,
out int offset)
{
offset = 0;
if (virtualAddress < baseAddress)
{
return false;
}
var relative = virtualAddress - baseAddress;
if (relative + (ulong)length > (ulong)_storage.Length)
{
return false;
}
offset = (int)relative;
return true;
}
}
}
@@ -1,197 +0,0 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Buffers.Binary;
using SharpEmu.ShaderCompiler;
using SharpEmu.ShaderCompiler.Vulkan;
using Xunit;
namespace SharpEmu.ShaderCompiler.Tests;
public sealed class Gen5ImageTests
{
private const ulong ShaderAddress = 0x1_0000_C000;
private const uint SEndpgm = 0xBF810000;
[Theory]
[InlineData(1u, SpirvImageDim.Dim2D, 2u)]
[InlineData(2u, SpirvImageDim.Dim3D, 3u)]
public void ImageStoreDimensionControlsImageAndCoordinateTypes(
uint dimension,
SpirvImageDim expectedImageDimension,
uint expectedCoordinateComponents)
{
var instructions = ReadSpirvInstructions(
CompileImageOperation("ImageStore", dimension));
var imageType = Assert.Single(
instructions,
item => item.Opcode == SpirvOp.TypeImage);
Assert.Equal((uint)expectedImageDimension, imageType.Operands[2]);
Assert.Equal(2u, imageType.Operands[6]);
AssertCoordinateVectorWidth(
instructions,
SpirvOp.ImageWrite,
coordinateOperand: 1,
expectedComponents: expectedCoordinateComponents);
var sizeQuery = Assert.Single(
instructions,
item => item.Opcode == SpirvOp.ImageQuerySize);
AssertVectorTypeWidth(
instructions,
sizeQuery.Operands[0],
expectedCoordinateComponents);
}
[Fact]
public void ImageSampleDim3DUsesThreeComponentSampleCoordinates()
{
var instructions = ReadSpirvInstructions(
CompileImageOperation("ImageSampleLz", dimension: 2));
var imageType = Assert.Single(
instructions,
item => item.Opcode == SpirvOp.TypeImage);
Assert.Equal((uint)SpirvImageDim.Dim3D, imageType.Operands[2]);
Assert.Equal(1u, imageType.Operands[6]);
AssertCoordinateVectorWidth(
instructions,
SpirvOp.ImageSampleExplicitLod,
coordinateOperand: 3,
expectedComponents: 3);
}
private static byte[] CompileImageOperation(string opcode, uint dimension)
{
var addressRegisters = dimension == 2
? new uint[] { 0, 1, 2 }
: [0, 1];
var control = new Gen5ImageControl(
Dmask: 0xF,
VectorAddress: 0,
AddressRegisters: addressRegisters,
VectorData: 4,
ScalarResource: 8,
ScalarSampler: 16,
Dimension: dimension,
IsArray: false,
Glc: false,
Slc: false,
A16: false,
D16: false);
var imageInstruction = new Gen5ShaderInstruction(
0,
Gen5ShaderEncoding.Mimg,
opcode,
[],
[],
[],
control);
var end = new Gen5ShaderInstruction(
8,
Gen5ShaderEncoding.Sopp,
"SEndpgm",
[SEndpgm],
[],
[],
null);
var state = new Gen5ShaderState(
new Gen5ShaderProgram(ShaderAddress, [imageInstruction, end]),
[],
null);
var scalarRegisters = new uint[256];
var descriptor = new uint[8];
descriptor[1] = 71u << 20; // FORMAT_16_16_16_16_FLOAT
descriptor[3] = (dimension == 2 ? 10u : 9u) << 28;
var evaluation = new Gen5ShaderEvaluation(
scalarRegisters,
scalarRegisters,
[
new Gen5ImageBinding(
imageInstruction.Pc,
imageInstruction.Opcode,
control,
descriptor,
new uint[4],
null),
],
[]);
Assert.True(
Gen5SpirvTranslator.TryCompileComputeShader(
state,
evaluation,
1,
1,
1,
out var shader,
out var error),
error);
return shader.Spirv;
}
private static IReadOnlyList<ParsedSpirvInstruction> ReadSpirvInstructions(
byte[] spirv)
{
var instructions = new List<ParsedSpirvInstruction>();
for (var offset = 5 * sizeof(uint); offset < spirv.Length;)
{
var instruction = BinaryPrimitives.ReadUInt32LittleEndian(
spirv.AsSpan(offset));
var wordCount = checked((int)(instruction >> 16));
Assert.InRange(wordCount, 1, (spirv.Length - offset) / sizeof(uint));
var operands = new uint[wordCount - 1];
for (var operand = 0; operand < operands.Length; operand++)
{
operands[operand] = BinaryPrimitives.ReadUInt32LittleEndian(
spirv.AsSpan(offset + (operand + 1) * sizeof(uint)));
}
instructions.Add(
new ParsedSpirvInstruction((SpirvOp)(ushort)instruction, operands));
offset += wordCount * sizeof(uint);
}
return instructions;
}
private static void AssertCoordinateVectorWidth(
IReadOnlyList<ParsedSpirvInstruction> instructions,
SpirvOp operation,
int coordinateOperand,
uint expectedComponents)
{
var imageOperation = Assert.Single(
instructions,
item => item.Opcode == operation);
var coordinateId = imageOperation.Operands[coordinateOperand];
var coordinate = Assert.Single(
instructions,
item =>
item.Opcode == SpirvOp.CompositeConstruct &&
item.Operands.Length >= 2 &&
item.Operands[1] == coordinateId);
AssertVectorTypeWidth(
instructions,
coordinate.Operands[0],
expectedComponents);
}
private static void AssertVectorTypeWidth(
IReadOnlyList<ParsedSpirvInstruction> instructions,
uint vectorTypeId,
uint expectedComponents)
{
var vectorType = Assert.Single(
instructions,
item =>
item.Opcode == SpirvOp.TypeVector &&
item.Operands[0] == vectorTypeId);
Assert.Equal(expectedComponents, vectorType.Operands[2]);
}
private readonly record struct ParsedSpirvInstruction(
SpirvOp Opcode,
uint[] Operands);
}
@@ -1,22 +0,0 @@
<!--
Copyright (C) 2026 SharpEmu Emulator Project
SPDX-License-Identifier: GPL-2.0-or-later
-->
<Project Sdk="Microsoft.NET.Sdk">
<PropertyGroup>
<IsPackable>false</IsPackable>
<GenerateDocumentationFile>false</GenerateDocumentationFile>
</PropertyGroup>
<ItemGroup>
<ProjectReference Include="..\..\src\SharpEmu.ShaderCompiler\SharpEmu.ShaderCompiler.csproj" />
<ProjectReference Include="..\..\src\SharpEmu.ShaderCompiler.Vulkan\SharpEmu.ShaderCompiler.Vulkan.csproj" />
</ItemGroup>
<ItemGroup>
<PackageReference Include="Microsoft.NET.Test.Sdk" />
<PackageReference Include="xunit" />
<PackageReference Include="xunit.runner.visualstudio" />
</ItemGroup>
</Project>