Compare commits

...

27 Commits

Author SHA1 Message Date
ParantezTech 2e5787de06 [CPU] Fix Sema ORBIS_GEN2_ERROR_BUSY loop 2026-07-25 14:35:28 +03:00
Andrew 8f9456229a fix(memory): reserve only large regions (#608)
* fix(memory): reserve only large regions

* Potential fix for pull request finding
2026-07-25 14:25:53 +03:00
Berk 5b602c0232 [GPU] Fix detiled cache key for VulkanDetilePass (#620) 2026-07-25 14:12:17 +03:00
Digote 26c502914c feat(audio): implement sceAudioOutOutputs (#605)
* feat(audio): implement batched output submission

* test(audio): cover batched output semantics

* test(audio): cover multi-port output batches

---------

Co-authored-by: diego <diego@DIGOTE-PC>
2026-07-24 20:13:30 +03:00
shadowbeat070 a158960c20 feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal) (#592)
* feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal)

Move RDNA2 exact-XOR deswizzle (swizzle modes 5/9/24/27, 4bpp) off the CPU
onto a GPU compute pass. GnmTiling.GetDetileParams resolves the shared
addressing into DetileParams; the CPU fallback and both GPU kernels consume
the same params so they never disagree.

Vulkan (verified bit-exact on NVIDIA): SpirvFixedShaders.CreateDetileCompute
hand-emits the SPIR-V kernel; VulkanDetilePass.RecordDetile records the
dispatch into the async batch command buffer (never a blocking submit on the
render thread) with transients retired via fence; VulkanDetileSelfTest
(SHARPEMU_DETILE_SELFTEST=1) checks both entry points against the CPU detile.

Metal (Mac-untested): detile_compute.msl (detile_cs) + MetalDetilePass mirror
the Vulkan pass. The active Metal path CPU-detiles via the new
GnmTiling.DetileWithParams when a texture arrives packaged (empty RgbaPixels +
TiledSource/Detile), keeping Metal correct under default-on with no regression;
wiring MetalDetilePass live is the remaining on-device step.

Flags: GPU detile is default-on (SHARPEMU_GPU_DETILE=0 disables);
[GPU-DETILE] diagnostics gated behind SHARPEMU_LOG_GPU_DETILE=1.

Tests: 17 detile unit tests pass, incl. DetileWithParams and GetDetileParams
each matching TryDetile bit-for-bit across all supported modes/bpp, plus a
SPIR-V structural-validity test.

* feat(gpu): GPU compute detile for guest tiled textures (Vulkan + Metal)

Move RDNA2 exact-XOR deswizzle (swizzle modes 5/9/24/27, 4bpp) off the CPU
onto a GPU compute pass. GnmTiling.GetDetileParams resolves the shared
addressing into DetileParams that the CPU fallback and both GPU kernels
consume, so they never disagree; everything else keeps the CPU path.

Vulkan (verified bit-exact on NVIDIA): SpirvFixedShaders.CreateDetileCompute
hand-emits the kernel; VulkanDetilePass.RecordDetile records into the async
batch command buffer (never a blocking submit on the render thread) with
transients retired via fence, falling back to CPU detile on failure.
VulkanDetileSelfTest (SHARPEMU_DETILE_SELFTEST=1) checks both entry points.

Metal (Mac-untested): detile_compute.msl + MetalDetilePass mirror the Vulkan
pass; the active Metal path CPU-detiles via GnmTiling.DetileWithParams so it
stays correct under default-on. Wiring MetalDetilePass live is a follow-up.

Flags: default-on (SHARPEMU_GPU_DETILE=0 disables); diagnostics behind
SHARPEMU_LOG_GPU_DETILE=1. Adds 17 passing detile unit tests.

* Fix: added support layered texture support for the GPU-Detiling.

* Fix: Added support for BlockTable (1 / 4 / 8 (Morton/Z-order))

* feat: added support for 8 and 16 bpp (bytes per element)

* Fixed a build failure specific to this branch

---------
2026-07-24 20:13:02 +03:00
MarcelMediaDev 5228335f15 fix(gpu): support Gen5 flat memory and 3D images (#587)
Vector-mesh UI text samples type-10 volume LUTs; treat MIMG DIM=2 as
Dim3D and transport depth through AGC and Vulkan so Z slices no longer
collapse into a single 2D plane.
2026-07-24 15:44:58 +03:00
Berk 21f964a0dc Update README 2026-07-24 03:00:05 +03:00
ParantezTech 6133313a83 [readme] added support SharpEmu section 2026-07-24 02:58:58 +03:00
ParantezTech 6db095ec82 revert: restore state before huge regression 2026-07-23 16:03:45 +03:00
CasualcoderDev 5a08a9bb43 fix: VulkanHostBufferPool deadlock, audio overflow crash, and log grouping (#564)
* fix: VulkanHostBufferPool deadlock, audio overflow crash, and log grouping

* fix: implement thread-safe buffer pool, refactor output handling, and use unchecked cast for audio conversion
2026-07-23 15:41:51 +03:00
MarcelMediaDev f9d92135a0 fix(agc): merge Prospero attrib-table formats onto IR vertex inputs (#556)
IR-discovered BufferLoadFormat often keeps a stale float sharp format;
patch DataFormat/offset from the AGC attrib table (semantic index),
allow offen fetches, and map quirks 113/121 through NarrowVk for host
vertex input.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:39:24 +03:00
Radu Ursache 8779c96c3a test(debugger): add C# unit tests for protocol and breakpoints (#568)
Cover DebugRequest parsing, BreakpointStore, and DebugCommandDispatcher
with a fake session. Wire the project into SharpEmu.slnx.
2026-07-23 15:38:08 +03:00
MarcelMediaDev 4c6cff1116 fix(agc): skip CB metadata draws for EliminateFastClear/Fmask/DCC (#553)
CB_COLOR_CONTROL modes 2/5/6 are colour-buffer metadata ops; applying
the bound shader as a normal colour draw corrupts subsequent composites.
Decode MODE from bits [6:4] and return before translate.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:37:34 +03:00
MarcelMediaDev 82ab181861 fix(hle): enable GuestImageWriteTracker CPU sync on Windows (#550)
Windows previously hard-disabled the tracker, so CPU-written guest
planes never marked dirty and host textures stayed empty. Arm pages
with VirtualProtect, handle write AVs in VEH, and warm/test on
VirtualAlloc memory so protect cannot poison the CRT heap.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:37:05 +03:00
MarcelMediaDev 007bf6fa73 fix(kernel): reject getdents on file fds and emit . / .. for empty dirs (#546)
Returning rax=0 for non-directory or empty listings looked like EOF and
let GTA treat the fd as a pointer (fiWriteAsyncDataWorker AV at 0xB1).

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:36:28 +03:00
MarcelMediaDev 8e1e89c024 fix(agc): accept Gen5 hull shaders that omit PGM_LO/HI in CreateShader (#545)
Type-5 headers can start with RSRC1/RSRC2; rejecting them left null handles
and Main Thread AVs. Scan the SH table and skip PGM patch when absent.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:36:19 +03:00
Kurt Himebauch 2764aaab3f feat: implement cosf, time, ctype tables, tracked heap access, IL2CPP lookup ABI paths (#542)
* fix: add Messenger CRT and AGC compatibility shims

* fix: keep Messenger IL2CPP bootstrap on HLE shims

* Fix Messenger compatibility ABI handling

* Make IL2CPP ABI regression tests portable
2026-07-23 15:35:59 +03:00
MarcelMediaDev 7b950166d7 fix(remoteplay): stub Initialize and GetConnectionStatus as disconnected (#536)
* fix(remoteplay): stub Initialize and GetConnectionStatus as disconnected

Titles probe Remote Play during pad/network bring-up; unresolved imports
returned NOT_FOUND. Report initialized + disconnected so callers take the
normal offline path.

Co-authored-by: Cursor <cursoragent@cursor.com>

* chore: retrigger gameplay CI for PR #536

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:35:28 +03:00
MarcelMediaDev eb1195e59a fix(kernel): implement APR ResolveFilepathsWithPrefixToIdsAndFileSizes (#534)
* fix(kernel): implement APR ResolveFilepathsWithPrefixToIdsAndFileSizes

Resource streamers resolve relative paths against a shared prefix; without
this HLE every call returned NOT_FOUND and assets never got real ids/sizes.

Co-authored-by: Cursor <cursoragent@cursor.com>

* chore: retrigger gameplay CI for PR #534

Co-authored-by: Cursor <cursoragent@cursor.com>

---------

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:34:58 +03:00
MarcelMediaDev e13cb28267 fix(audio): harden AudioOut2 stack out-buffer writes against canary smash (#532)
Titles that stack-allocate AudioOut2 outs next to the frame canary were
corrupted by oversized or mistyped HLE writes; keep ContextPush pacing.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 15:34:40 +03:00
Radu Ursache 956da769a3 fix(kernel): finish Posix -1/errno for file ops and open EACCES (#567)
Map UnauthorizedAccess on open to PERMISSION_DENIED and route Posix
lseek/pread/pwrite/rename/etc failures through PosixFailure so libc-style
callers get RAX=-1 plus TLS errno, matching open/read/write.
2026-07-23 15:32:56 +03:00
h4sht d7bd814fb9 [Core/Dlsym] Restore sceKernelDlsym bootstrap argument normalization from PR #94 (#565)
* [Core/Dlsym] Restore normalize-dlsym-arguments and deferred bootstrap tracing from PR #94

PR #216 regressed two critical features from PR #94:
1. NormalizeKernelDynlibDlsymArguments — handles argument reordering
   when standalone bootstrap loaders call sceKernelDlsym through the
   bridge with (symbol_ptr, handle, out) instead of the standard
   (handle, symbol_ptr, out). Without this, payloads like elfldr-ps5
   and websrv-ps5 fail with a deterministic UnmanagedCallersOnly
   fail-fast.
2. Deferred bootstrap tracing — ring-buffered import logging that
   drains after the hot path, avoiding per-call Console.Error I/O.

Also restored:
- CompleteKernelDynlibDlsymFailure — centralized error handling
- IsPlausibleDynlibSymbolPointer — pointer bounds validation
- COW snapshot of _importEntries in ProbeReturnRip
- ResetLazyDlsymStubState and lazy-dlsym field infrastructure
- DraftDrainDeferredBootstrapTraces in Execute() finally block

Fixes #530, fixes #531

* fix: remove orphaned _importNidHashCache.Clear() reference

The field _importNidHashCache no longer exists on main (removed post PR #94).
The 3-way merge incorrectly restored the .Clear() call without the field
declaration, causing a build failure on all platforms.

---------

Co-authored-by: tru3 <tru3@tru3.com>
2026-07-23 15:24:59 +03:00
Mike Saito 96fde5764f Astro Bot stack: VEH/TBB, title clear, swapchain fallback, Psml MFSR (#528)
* Cpu/Kernel: harden VEH trampoline and keep TBB on native workers

Route FastFail/CLR/stack-overflow around managed VEH, serialize managed
entry with a recursive spinlock, require native workers for tbb_thead,
and abandon pthread mutexes when a guest thread is torn down by worker
abort so splash waiters are not left holding locks forever.

* Cpu: soft-fail TBB native worker storms and cap concurrent Runs

Throwing on worker/prologue faults killed the process mid tbb_thead
burst (FailFast 0xC0000409). Soft-return 0x80020012, limit in-flight
native Runs (default 2), and keep prewarm small so back-to-back boots
do not need an artificial settle delay.

* Agc/VideoOut: poison-only empty-SRT reject and clear procedural ES/PS

Skip QueueSubmit only when Address-0 image slots remain; run the
Astro title clear pair via CmdClearColorImage so the pass executes
without descriptors that lose the device.

* VideoOut: recreate swapchain with fallback extent on 0x0 surface

Minimized Win32 surfaces report 0x0 / MaxImageExtent=0; deferring
recreate forever left an OutOfDate swapchain with no presents.
Clamp to last/default size and recreate instead of early-return.

* Psml: stub MFSR init/shared/context and dispatch packet size

Astro Bot asserts in GfxRenderStagePSSR when scePsmlMfsrInit is unresolved
(Mfsr initialized failed). Soft HLE for the MFSR shared-resource and
800M3_2 context path plus dispatch packet size lets boot pass splash to
first frame without claiming real upscaling.

* Psml: stub MFSR GetDispatchMfsrPacket900 for logo PSSR

Astro StartLevel ps_logo asserted GfxRenderStagePSSR.cpp:266 when
GetDispatchMfsrPacket900 (RUNLFro+qok) was unresolved. Return SCE_OK
from SizeInDwords so the 900 fill runs, soft-clear the guest packet
buffer, and register 1000/1100 siblings for the same ABI.
2026-07-23 12:37:44 +03:00
MarcelMediaDev 7a108c6f87 fix(hle): GameService Ok stubs plus NetInetPton and Json terminate (#560)
Resolve logo-accept-gate unresolved NIDs seen on PPSA04264: share/voice/
telemetry/content Ok stubs, sceNetInetPton, and Json Initializer::terminate.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 12:33:39 +03:00
Kurt Himebauch eb252af7b3 unstub: preserve notice screen skip flag (#559)
* fix(system-service): preserve notice screen skip flag

* test(bink): avoid frame timing flake on CI
2026-07-23 12:32:49 +03:00
MarcelMediaDev 8e5a0bfb19 fix(agc): implement sceAgcDcbSetUcRegisterDirect (#558)
Emit the 3-dword SET_UCONFIG_REG packet from the packed {offset,value}
in RSI. Unresolved calls left GPU config registers unset during
RenderThread/Main bring-up.

Co-authored-by: Cursor <cursoragent@cursor.com>
2026-07-23 12:32:04 +03:00
ParantezTech d991e32b15 [readme] update screenshots 2026-07-23 03:14:26 +03:00
47 changed files with 7424 additions and 545 deletions
Binary file not shown.

Before

Width:  |  Height:  |  Size: 190 KiB

After

Width:  |  Height:  |  Size: 345 KiB

Binary file not shown.

Before

Width:  |  Height:  |  Size: 101 KiB

After

Width:  |  Height:  |  Size: 227 KiB

Binary file not shown.

Before

Width:  |  Height:  |  Size: 86 KiB

After

Width:  |  Height:  |  Size: 104 KiB

+20
View File
@@ -25,6 +25,14 @@ 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
@@ -136,6 +144,18 @@ 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,6 +21,7 @@ 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,6 +153133,7 @@ scePsmlMfsrGetContextBufferRequirement800M3_2
scePsmlMfsrGetDispatchMfsrPacket1000
scePsmlMfsrGetDispatchMfsrPacket1100
scePsmlMfsrGetDispatchMfsrPacketSizeInDwords
scePsmlMfsrGetDispatchMfsrPacket900
scePsmlMfsrGetMipmapBias
scePsmlMfsrGetSharedResourcesInitRequirement
scePsmlMfsrInit
@@ -19,6 +19,9 @@ 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()
{
@@ -435,18 +438,91 @@ public sealed partial class DirectExecutionBackend
void* contextRecord,
ulong rip)
{
if (exceptionRecord->ExceptionCode != 3221225477u ||
rip >= 0x0000000800000000UL ||
_activeGuestThreadState is not { Name: "tbb_thead" } activeThread)
if (exceptionRecord->ExceptionCode != 3221225477u)
{
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(
@@ -458,13 +534,57 @@ public sealed partial class DirectExecutionBackend
_ = TryPatchActiveGuestReturnSlot(hostExit);
WriteCtxU64(contextRecord, 120, 0);
WriteCtxU64(contextRecord, 248, hostExit);
var recovery = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultRecoveries);
var recoveryFallback = Interlocked.Increment(ref _auxiliaryThreadExecuteFaultRecoveries);
Console.Error.WriteLine(
$"[LOADER][WARN] Recovered auxiliary TBB execute fault #{recovery}: " +
$"[LOADER][WARN] Recovered auxiliary TBB execute fault #{recoveryFallback}: " +
$"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,6 +1449,9 @@ 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);
@@ -1463,6 +1466,7 @@ public sealed partial class DirectExecutionBackend
!expectedEqueueTimeout &&
!expectedMutexTrylockBusy &&
!expectedSemaphoreTrywaitAgain &&
!expectedPollSemaBusy &&
!expectedNetAcceptWouldBlock &&
!expectedUserServiceNoEvent &&
!expectedPrivacyInvalidParameter)
@@ -29,9 +29,29 @@ 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;
@@ -49,6 +69,9 @@ 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.
@@ -56,40 +79,148 @@ 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)
private unsafe int RunGuestEntryStub(void* entryStub, ulong hostRspSlot, bool requireNativeWorker = false)
{
var worker = RentNativeGuestExecutor();
if (worker is null)
{
TlsSetValue(_hostRspSlotTlsIndex, (nint)hostRspSlot);
return CallNativeEntry(entryStub);
}
// 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;
try
{
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;
// 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);
}
}
finally
{
ReturnNativeGuestExecutor(worker);
_nativeWorkerRunLimiter.Release();
}
}
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
@@ -400,6 +531,22 @@ 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,
@@ -445,6 +592,49 @@ 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;
@@ -452,7 +642,22 @@ public sealed partial class DirectExecutionBackend
forcedExit = _runForcedExit;
if (_runPrologueFailed)
{
throw new InvalidOperationException("Native guest worker failed to bind the run ambient (prologue fault)");
// 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);
}
return _runNativeResult;
}
@@ -546,6 +751,18 @@ 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);
@@ -580,6 +797,14 @@ 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,6 +214,15 @@ 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;
@@ -1041,6 +1050,8 @@ 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");
@@ -1064,6 +1075,7 @@ 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
{
@@ -1087,13 +1099,29 @@ 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)
@@ -2409,9 +2437,147 @@ 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)
{
const uint stubSize = 256u;
// 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;
void* ptr = VirtualAlloc(null, stubSize, 12288u, 64u);
if (ptr == null)
{
@@ -2420,10 +2586,305 @@ 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);
@@ -2440,11 +2901,67 @@ 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;
@@ -2470,11 +2987,64 @@ 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;
@@ -2495,6 +3065,18 @@ 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);
@@ -5149,16 +5731,27 @@ 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
{
var nativeReturn = CallNativeEntry(ptr);
// 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);
}
if (ActiveGuestThreadYieldRequested)
{
reason = ActiveGuestThreadYieldReason ?? "guest thread blocked";
@@ -5304,16 +5897,24 @@ 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
{
var nativeReturn = CallNativeEntry(ptr);
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);
}
if (ActiveGuestThreadYieldRequested)
{
reason = ActiveGuestThreadYieldReason ?? "guest thread blocked";
@@ -6446,6 +7047,16 @@ 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);
@@ -6456,6 +7067,11 @@ 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);
@@ -6466,6 +7082,11 @@ 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 = 256u;
const uint stubSize = 1024u;
void* ptr = (void*)_memory.Allocate(0, stubSize, HostPageProtection.ReadWriteExecute);
if (ptr == null)
{
@@ -43,11 +43,15 @@ 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
@@ -55,13 +59,162 @@ 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, 0xEB); EmitByte(code, ref offset, 0x03); // jmp over pass block
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); // 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,7 +238,15 @@ 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;
var result = _hostMemory.Allocate(desiredAddress, alignedSize, hostProtection);
// 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);
if (result == 0)
{
return false;
@@ -252,6 +260,8 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
return false;
}
var state = reservedOnly ? ReserveRegion(actualAddress, alignedSize) : "n/a";
_gate.EnterWriteLock();
try
{
@@ -260,7 +270,7 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
VirtualAddress = actualAddress,
Size = alignedSize,
IsExecutable = executable,
IsReservedOnly = false,
IsReservedOnly = reservedOnly,
Protection = protection
});
}
@@ -269,6 +279,7 @@ 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;
@@ -361,44 +372,7 @@ public sealed unsafe class PhysicalVirtualMemory : IVirtualMemory, IGuestMemoryA
var actualAddress = result;
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";
}
}
var lazyPrimeState = reservedOnly ? ReserveRegion(actualAddress, alignedSize) : "n/a";
_gate.EnterWriteLock();
try
@@ -425,6 +399,41 @@ 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,6 +221,29 @@ 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,6 +6,55 @@ 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
@@ -501,6 +550,141 @@ 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,
+279 -14
View File
@@ -4,6 +4,7 @@
using SharpEmu.HLE;
using SharpEmu.HLE.Host;
using System.Buffers;
using System.Buffers.Binary;
using System.Collections.Concurrent;
using System.Diagnostics;
@@ -11,8 +12,17 @@ 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.
@@ -56,6 +66,7 @@ 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);
@@ -83,7 +94,24 @@ public static class AudioOutExports
}
}
public void Dispose() => Backend?.Dispose();
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;
}
[SysAbiExport(
@@ -115,9 +143,18 @@ public static class AudioOutExports
string backendName;
try
{
var audio = HostPlatform.Current.Audio;
backend = audio.OpenStereoPcm16Stream(frequency);
backendName = audio.BackendName;
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;
}
}
catch (Exception exception)
{
@@ -192,6 +229,46 @@ 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",
@@ -225,16 +302,7 @@ public static class AudioOutExports
return ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
}
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")}");
}
}
TraceOutput(handle, port, source);
if (port.Backend is null)
{
@@ -272,6 +340,184 @@ 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",
@@ -362,6 +608,25 @@ 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(
+16 -3
View File
@@ -1,6 +1,8 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using SharpEmu.Libs.Agc;
namespace SharpEmu.Libs.Gpu;
// The types that cross the guest-GPU backend seam. Every field is either a neutral
@@ -10,7 +12,8 @@ namespace SharpEmu.Libs.Gpu;
// translation for its API.
/// <summary>A guest texture referenced by a draw or dispatch. Format/NumberType/
/// TileMode/DstSelect are raw guest descriptor codes.</summary>
/// TileMode/DstSelect/Type are raw guest descriptor codes. Depth is the
/// normalized volume depth (one for non-3D resources).</summary>
internal sealed record GuestDrawTexture(
ulong Address,
uint Width,
@@ -32,7 +35,15 @@ 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 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);
/// <summary>Raw guest sampler descriptor dwords, copied verbatim from guest memory.</summary>
internal readonly record struct GuestSampler(
@@ -55,7 +66,9 @@ internal readonly record struct TextureContentIdentity(
uint Pitch,
GuestSampler Sampler,
bool Arrayed = false,
uint ArrayLayers = 1);
uint ArrayLayers = 1,
uint Type = 9,
uint Depth = 1);
internal sealed record GuestMemoryBuffer(
ulong BaseAddress,
@@ -0,0 +1,313 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Numerics;
using SharpEmu.Libs.Agc;
using SharpEmu.ShaderCompiler.Metal;
namespace SharpEmu.Libs.Gpu.Metal;
/// <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,6 +1,7 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using SharpEmu.Libs.Agc;
using SharpEmu.ShaderCompiler;
using SharpEmu.ShaderCompiler.Metal;
@@ -1583,6 +1584,45 @@ internal static partial class MetalVideoPresenter
return cached;
}
// Default-on GPU detile packages the tiled source + resolved DetileParams
// with empty RgbaPixels so a backend can deswizzle on the GPU. The Metal
// GPU compute pass (MetalDetilePass / detile_compute.msl) is the intended
// equivalent of VulkanDetilePass, but it is Mac-untested, so the active
// Metal path CPU-detiles here via GnmTiling.DetileWithParams — the exact
// same DetileParams addressing the kernel runs — restoring the linear-upload
// behavior Metal had before GPU detile existed, with no regression.
if (texture.RgbaPixels.Length == 0 &&
texture.TiledSource is { } tiledSource &&
texture.Detile is { } detileParameters)
{
// The tiled source packs the array slices contiguously (one per layer);
// detile each into its layer-major linear region so the reconstructed
// pixels match what the CPU array-upload path produced pre-GPU-detile.
// A plain 2D texture is just one layer.
var layers = Math.Max((int)texture.ArrayLayers, 1);
var sliceLinearBytes =
detileParameters.ElementsWide * detileParameters.ElementsHigh * detileParameters.BytesPerElement;
var sliceTiledBytes = tiledSource.Length / layers;
var linear = new byte[sliceLinearBytes * layers];
var detiledAll = true;
for (var layer = 0; layer < layers; layer++)
{
if (!GnmTiling.DetileWithParams(
detileParameters,
tiledSource.AsSpan(layer * sliceTiledBytes, sliceTiledBytes),
linear.AsSpan(layer * sliceLinearBytes, sliceLinearBytes)))
{
detiledAll = false;
break;
}
}
if (detiledAll)
{
texture = texture with { RgbaPixels = linear };
}
}
// AGC ships the raw (detiled) source texels; create the texture in the
// guest's native format — Mac-family GPUs sample BC blocks directly —
// and size expectations with the same block-aware math AGC used.
+12
View File
@@ -80,6 +80,18 @@ 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,6 +153,64 @@ 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(
+31 -3
View File
@@ -27,8 +27,12 @@ internal static class KernelPthreadState
internal static ulong GetCurrentThreadHandle()
{
var guestThreadHandle = GuestThreadExecution.CurrentGuestThreadHandle;
if (guestThreadHandle != 0 && TryGetThreadIdentity(guestThreadHandle, out _))
// 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)
{
EnsureGuestThreadIdentity(guestThreadHandle);
return guestThreadHandle;
}
@@ -39,15 +43,27 @@ internal static class KernelPthreadState
internal static ulong GetCurrentThreadUniqueId()
{
var guestThreadHandle = GuestThreadExecution.CurrentGuestThreadHandle;
if (guestThreadHandle != 0 && TryGetThreadIdentity(guestThreadHandle, out var identity))
if (guestThreadHandle != 0)
{
return identity.UniqueId;
return EnsureGuestThreadIdentity(guestThreadHandle).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));
@@ -59,6 +75,18 @@ 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,6 +755,52 @@ 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
@@ -0,0 +1,365 @@
// 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,4 +59,65 @@ 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,6 +18,7 @@ public static class SystemServiceExports
private const int TitleIdFieldSize = 0x10;
private static string? _mainAppTitleId;
private static int _noticeScreenSkipFlag;
public static void ConfigureApplicationInfo(string? titleId)
{
@@ -37,9 +38,11 @@ public static class SystemServiceExports
return ctx.SetReturn(OrbisSystemServiceErrorParameter);
}
// No system notice screen to skip in the emulator; report "do not skip".
// 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.
Span<byte> flagBytes = stackalloc byte[1];
flagBytes[0] = 0;
flagBytes[0] = unchecked((byte)Volatile.Read(ref _noticeScreenSkipFlag));
return ctx.Memory.TryWrite(flagAddress, flagBytes)
? ctx.SetReturn(0)
: ctx.SetReturn((int)OrbisGen2Result.ORBIS_GEN2_ERROR_MEMORY_FAULT);
@@ -61,8 +64,13 @@ public static class SystemServiceExports
ExportName = "sceSystemServiceSetNoticeScreenSkipFlag",
Target = Generation.Gen5,
LibraryName = "libSceSystemService")]
public static int SystemServiceSetNoticeScreenSkipFlag(CpuContext ctx) =>
ctx.SetReturn(0);
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",
@@ -231,4 +239,7 @@ 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);
}
@@ -0,0 +1,894 @@
// 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);
}
}
}
@@ -0,0 +1,482 @@
// 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
@@ -76,6 +76,14 @@ 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";
}
@@ -0,0 +1,70 @@
#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,7 +314,8 @@ public static partial class Gen5SpirvTranslator
uint VectorType,
ImageComponentKind ComponentKind,
bool IsStorage,
bool Arrayed);
bool Arrayed,
SpirvImageDim Dimension);
private readonly record struct SpirvVertexInput(
uint Variable,
@@ -1001,11 +1002,15 @@ public static partial class Gen5SpirvTranslator
SpirvCapability.StorageImageExtendedFormats);
}
var isArrayed = !isStorage &&
var dimension = binding.Control.Dimension == 2
? SpirvImageDim.Dim3D
: SpirvImageDim.Dim2D;
var isArrayed = dimension != SpirvImageDim.Dim3D &&
!isStorage &&
Gen5ShaderTranslator.IsArrayedImageBinding(binding);
var imageType = _module.TypeImage(
componentType,
SpirvImageDim.Dim2D,
dimension,
depth: false,
arrayed: isArrayed,
multisampled: false,
@@ -1035,7 +1040,8 @@ public static partial class Gen5SpirvTranslator
_module.TypeVector(componentType, 4),
componentKind,
isStorage,
isArrayed));
isArrayed,
dimension));
_interfaces.Add(variable);
}
}
@@ -2259,20 +2265,39 @@ 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(
LoadV(control.VectorAddress),
vectorByteAddress,
UInt(unchecked((uint)control.OffsetBytes)));
byteAddress = ApplyGuestBufferByteBias(bindingIndex, byteAddress);
var dwordAddress = ShiftRightLogical(byteAddress, UInt(2));
if (instruction.Opcode is "GlobalAtomicAdd" or "GlobalAtomicUMax")
if (memoryOpcode is "GlobalAtomicAdd" or "GlobalAtomicUMax")
{
EmitExecConditional(() =>
{
EmitConditional(IsBufferWordInRange(bindingIndex, dwordAddress), () =>
{
var original = _module.AddInstruction(
instruction.Opcode == "GlobalAtomicAdd"
memoryOpcode == "GlobalAtomicAdd"
? SpirvOp.AtomicIAdd
: SpirvOp.AtomicUMax,
_uintType,
@@ -2289,12 +2314,12 @@ public static partial class Gen5SpirvTranslator
return true;
}
if (instruction.Opcode.StartsWith("GlobalStore", StringComparison.Ordinal))
if (memoryOpcode.StartsWith("GlobalStore", StringComparison.Ordinal))
{
EmitExecConditional(() =>
{
if (TryGetSubdwordStoreInfo(
instruction.Opcode,
memoryOpcode,
out var byteCount,
out var sourceShift))
{
@@ -2324,7 +2349,7 @@ public static partial class Gen5SpirvTranslator
}
if (TryGetSubdwordLoadInfo(
instruction.Opcode,
memoryOpcode,
out var loadByteCount,
out var signExtend,
out var d16,
@@ -3268,6 +3293,7 @@ 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(
@@ -3278,7 +3304,7 @@ public static partial class Gen5SpirvTranslator
resource.IsStorage
? SpirvOp.ImageQuerySize
: SpirvOp.ImageQuerySizeLod,
_module.TypeVector(_intType, 2),
_module.TypeVector(_intType, sizeComponentCount),
resource.IsStorage
? [queryImage]
: [queryImage, UInt(0)]);
@@ -3291,7 +3317,7 @@ public static partial class Gen5SpirvTranslator
}
uint value;
if (component < 2)
if (component < sizeComponentCount)
{
var signedValue = _module.AddInstruction(
SpirvOp.CompositeExtract,
@@ -3319,7 +3345,12 @@ public static partial class Gen5SpirvTranslator
return false;
}
var coordinates = BuildIntegerCoordinates(image, 0);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var coordinates = BuildIntegerCoordinates(
image,
0,
coordinateComponentCount);
var components = new uint[4];
uint sourceIndex = 0;
for (var component = 0; component < components.Length; component++)
@@ -3355,13 +3386,14 @@ public static partial class Gen5SpirvTranslator
components);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, 2),
_module.TypeVector(_intType, coordinateComponentCount),
imageObject);
EmitBoundsCheckedImageWrite(
coordinates,
imageSize,
imageObject,
texel);
texel,
coordinateComponentCount);
return true;
}
@@ -3382,14 +3414,17 @@ public static partial class Gen5SpirvTranslator
}
var signed = resource.ComponentKind == ImageComponentKind.Sint;
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var atomicImageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, 2),
_module.TypeVector(_intType, coordinateComponentCount),
imageObject);
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
atomicImageSize);
atomicImageSize,
coordinateComponentCount);
EmitExecConditional(() =>
{
var pointer = _module.AddInstruction(
@@ -3433,14 +3468,17 @@ public static partial class Gen5SpirvTranslator
{
if (resource.IsStorage)
{
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySize,
_module.TypeVector(_intType, 2),
_module.TypeVector(_intType, coordinateComponentCount),
imageObject);
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
imageSize);
imageSize,
coordinateComponentCount);
sampled = _module.AddInstruction(
SpirvOp.ImageRead,
resource.VectorType,
@@ -3454,15 +3492,18 @@ public static partial class Gen5SpirvTranslator
SpirvOp.Image,
resource.ImageType,
imageObject);
var coordinateComponentCount =
ImageCoordinateComponentCount(resource);
var imageSize = _module.AddInstruction(
SpirvOp.ImageQuerySizeLod,
_module.TypeVector(_intType, 2),
_module.TypeVector(_intType, coordinateComponentCount),
fetchedImage,
UInt(mipLevel));
var coordinates = BuildClampedIntegerCoordinates(
image,
0,
imageSize);
imageSize,
coordinateComponentCount);
sampled = _module.AddInstruction(
SpirvOp.ImageFetch,
resource.VectorType,
@@ -3494,12 +3535,19 @@ 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);
offset = BuildImageOffset(
image,
addressCursor,
spatialComponentCount);
addressCursor += ImageFullAddressSlots(image);
}
@@ -3523,26 +3571,33 @@ public static partial class Gen5SpirvTranslator
}
var gradientX = hasGradients
? BuildFloatCoordinates(image, addressCursor)
? BuildFloatCoordinates(
image,
addressCursor,
spatialComponentCount)
: 0u;
var gradientY = hasGradients
? BuildFloatCoordinates(image, addressCursor + 2)
? BuildFloatCoordinates(
image,
addressCursor + (int)spatialComponentCount,
spatialComponentCount)
: 0u;
if (hasGradients)
{
addressCursor += 4;
addressCursor += checked((int)(spatialComponentCount * 2));
}
var coordinates = resource.Arrayed
? BuildFloatArrayCoordinates(image, addressCursor)
: BuildFloatCoordinates(image, addressCursor);
var coordinates = BuildFloatCoordinates(
image,
addressCursor,
coordinateComponentCount);
var explicitLod = hasGradients || hasZeroLod || hasLod;
var lod = hasZeroLod
? Float(0)
: hasLod
? LoadImageFloatAddress(
image,
addressCursor + (resource.Arrayed ? 3 : 2))
addressCursor + (int)coordinateComponentCount)
: lodOrBias;
if (hasOffset)
{
@@ -3608,11 +3663,18 @@ 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);
offset = BuildImageOffset(
image,
addressCursor,
spatialComponentCount);
addressCursor += ImageFullAddressSlots(image);
}
@@ -3627,9 +3689,10 @@ public static partial class Gen5SpirvTranslator
addressCursor += ImageFullAddressSlots(image);
}
var coordinates = resource.Arrayed
? BuildFloatArrayCoordinates(image, addressCursor)
: BuildFloatCoordinates(image, addressCursor);
var coordinates = BuildFloatCoordinates(
image,
addressCursor,
coordinateComponentCount);
var operands = new List<uint>
{
imageObject,
@@ -3822,28 +3885,31 @@ public static partial class Gen5SpirvTranslator
});
}
private uint BuildFloatCoordinates(Gen5ImageControl image, int start)
{
var x = LoadImageFloatAddress(image, start);
var y = LoadImageFloatAddress(image, start + 1);
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_vec2Type,
x,
y);
}
private static uint ImageSpatialComponentCount(
SpirvImageResource resource) =>
resource.Dimension == SpirvImageDim.Dim3D ? 3u : 2u;
private uint BuildFloatArrayCoordinates(Gen5ImageControl image, int start)
private static uint ImageCoordinateComponentCount(
SpirvImageResource resource) =>
resource.Arrayed ? 3u : ImageSpatialComponentCount(resource);
private uint BuildFloatCoordinates(
Gen5ImageControl image,
int start,
uint componentCount)
{
var x = LoadImageFloatAddress(image, start);
var y = LoadImageFloatAddress(image, start + 1);
var slice = LoadImageFloatAddress(image, start + 2);
var components = new uint[checked((int)componentCount)];
for (var component = 0; component < components.Length; component++)
{
components[component] = LoadImageFloatAddress(
image,
start + component);
}
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
_vec3Type,
x,
y,
slice);
_module.TypeVector(_floatType, componentCount),
components);
}
private static int ImageAddressRegister(
@@ -3948,47 +4014,49 @@ public static partial class Gen5SpirvTranslator
ShiftLeftLogical(BitwiseAnd(high, UInt(0xFFFF)), UInt(16)));
}
private uint BuildIntegerCoordinates(Gen5ImageControl image, int start)
private uint BuildIntegerCoordinates(
Gen5ImageControl image,
int start,
uint componentCount)
{
var ivec2 = _module.TypeVector(_intType, 2);
var x = Bitcast(_intType, LoadImageIntegerAddress(image, start));
var y = Bitcast(_intType, LoadImageIntegerAddress(image, start + 1));
var components = new uint[checked((int)componentCount)];
for (var component = 0; component < components.Length; component++)
{
components[component] = Bitcast(
_intType,
LoadImageIntegerAddress(image, start + component));
}
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
ivec2,
x,
y);
_module.TypeVector(_intType, componentCount),
components);
}
private uint BuildClampedIntegerCoordinates(
Gen5ImageControl image,
int start,
uint imageSize)
uint imageSize,
uint componentCount)
{
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));
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));
}
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
ivec2,
x,
y);
_module.TypeVector(_intType, componentCount),
components);
}
private uint ClampSignedCoordinate(uint value, uint extent)
@@ -4027,69 +4095,47 @@ public static partial class Gen5SpirvTranslator
uint coordinates,
uint imageSize,
uint imageObject,
uint texel)
uint texel,
uint coordinateComponentCount)
{
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 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 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 writeLabel = _module.AllocateId();
var mergeLabel = _module.AllocateId();
_module.AddStatement(SpirvOp.SelectionMerge, mergeLabel, 0);
@@ -4108,30 +4154,30 @@ public static partial class Gen5SpirvTranslator
_module.AddLabel(mergeLabel);
}
private uint BuildImageOffset(Gen5ImageControl image, int component)
private uint BuildImageOffset(
Gen5ImageControl image,
int component,
uint componentCount)
{
var ivec2 = _module.TypeVector(_intType, 2);
var packed = Bitcast(
_intType,
LoadV(image.GetAddressRegister(
ImageAddressRegister(image, component))));
var x = _module.AddInstruction(
SpirvOp.BitFieldSExtract,
_intType,
packed,
UInt(0),
UInt(6));
var y = _module.AddInstruction(
SpirvOp.BitFieldSExtract,
_intType,
packed,
UInt(8),
UInt(6));
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));
}
return _module.AddInstruction(
SpirvOp.CompositeConstruct,
ivec2,
x,
y);
_module.TypeVector(_intType, componentCount),
components);
}
private uint ApplyDynamicSampleOffset(
@@ -4141,7 +4187,16 @@ public static partial class Gen5SpirvTranslator
uint texelOffset,
uint lod)
{
var ivec2 = _module.TypeVector(_intType, 2);
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 image = _module.AddInstruction(
SpirvOp.Image,
resource.ImageType,
@@ -4163,14 +4218,14 @@ public static partial class Gen5SpirvTranslator
signedLod);
var size = _module.AddInstruction(
SpirvOp.ImageQuerySizeLod,
resource.Arrayed ? _module.TypeVector(_intType, 3) : ivec2,
_module.TypeVector(_intType, queryComponentCount),
image,
clampedLod);
if (resource.Arrayed)
{
size = _module.AddInstruction(
SpirvOp.VectorShuffle,
ivec2,
spatialIntegerType,
size,
size,
0u,
@@ -4179,27 +4234,27 @@ public static partial class Gen5SpirvTranslator
var sizeFloat = _module.AddInstruction(
SpirvOp.ConvertSToF,
_vec2Type,
spatialFloatType,
size);
var offsetFloat = _module.AddInstruction(
SpirvOp.ConvertSToF,
_vec2Type,
spatialFloatType,
texelOffset);
var normalizedOffset = _module.AddInstruction(
SpirvOp.FDiv,
_vec2Type,
spatialFloatType,
offsetFloat,
sizeFloat);
if (!resource.Arrayed)
{
return _module.AddInstruction(
SpirvOp.FAdd,
_vec2Type,
spatialFloatType,
coordinates,
normalizedOffset);
}
var offsetVec3 = _module.AddInstruction(
var arrayOffset = _module.AddInstruction(
SpirvOp.CompositeConstruct,
_vec3Type,
_module.AddInstruction(
@@ -4217,7 +4272,7 @@ public static partial class Gen5SpirvTranslator
SpirvOp.FAdd,
_vec3Type,
coordinates,
offsetVec3);
arrayOffset);
}
private bool TryEmitExport(
@@ -260,4 +260,233 @@ 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();
}
}
+2 -1
View File
@@ -192,7 +192,8 @@ public sealed record Gen5GlobalMemoryControl(
uint ScalarAddress,
int OffsetBytes,
bool Glc,
bool Slc) : Gen5InstructionControl;
bool Slc,
bool UsesFlatAddress = false) : Gen5InstructionControl;
public sealed record Gen5BufferMemoryControl(
uint DwordCount,
@@ -4,6 +4,7 @@
using SharpEmu.HLE;
using System.Buffers;
using System.Buffers.Binary;
using System.Collections.Concurrent;
using System.Diagnostics;
using System.Numerics;
@@ -35,6 +36,14 @@ 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
@@ -349,7 +358,10 @@ public static class Gen5ShaderScalarEvaluator
if (globalMemory.ScalarAddress >= ScalarRegisterCount - 1)
{
error =
$"global-address-register-range pc=0x{instruction.Pc:X} " +
$"{(globalMemory.UsesFlatAddress
? "flat-address-base-unresolved"
: "global-address-register-range")} " +
$"pc=0x{instruction.Pc:X} " +
$"s{globalMemory.ScalarAddress}";
return false;
}
@@ -364,11 +376,18 @@ 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))
{
@@ -2101,6 +2120,15 @@ 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,6 +305,23 @@ 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
@@ -454,7 +471,15 @@ public static class Gen5ShaderTranslator
}
}
instructions.Add(CreateInstruction(pc, encoding, name, words));
var instruction = CreateInstruction(pc, encoding, name, words);
if (instruction.Control is Gen5GlobalMemoryControl
{
UsesFlatAddress: true,
})
{
instruction = ResolveFlatAddressBase(instructions, instruction);
}
instructions.Add(instruction);
instructionCount++;
pc += sizeDwords * sizeof(uint);
@@ -1378,35 +1403,42 @@ public static class Gen5ShaderTranslator
var opcode = (word >> 18) & 0x7F;
sizeDwords = 2;
error = string.Empty;
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,
}
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
: string.Empty;
return FinishDecode(
@@ -1629,6 +1661,82 @@ 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,
@@ -2040,7 +2148,13 @@ public static class Gen5ShaderTranslator
var vectorAddress = extra & 0xFF;
var vectorData = (extra >> 8) & 0xFF;
var scalarAddress = (extra >> 16) & 0x7F;
var dwordCount = opcode switch
var usesFlatAddress = opcode.StartsWith(
"Flat",
StringComparison.Ordinal);
var memoryOpcode = usesFlatAddress
? "Global" + opcode["Flat".Length..]
: opcode;
var dwordCount = memoryOpcode switch
{
"GlobalLoadUbyte" or
"GlobalLoadSbyte" or
@@ -2068,12 +2182,20 @@ public static class Gen5ShaderTranslator
"GlobalStoreDwordx4" => 4u,
_ => 0u,
};
sources =
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Scalar(scalarAddress),
];
destinations = opcode.StartsWith("GlobalLoad", StringComparison.Ordinal)
sources = usesFlatAddress
?
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Vector(vectorAddress + 1),
]
:
[
Gen5Operand.Vector(vectorAddress),
Gen5Operand.Scalar(scalarAddress),
];
destinations = memoryOpcode.StartsWith(
"GlobalLoad",
StringComparison.Ordinal)
? Enumerable
.Range((int)vectorData, checked((int)dwordCount))
.Select(index => Gen5Operand.Vector((uint)index))
@@ -2083,10 +2205,11 @@ public static class Gen5ShaderTranslator
dwordCount,
vectorAddress,
vectorData,
scalarAddress,
usesFlatAddress ? uint.MaxValue : scalarAddress,
SignExtend(word & 0x1FFF, 13),
((word >> 16) & 1) != 0,
((word >> 17) & 1) != 0);
((word >> 17) & 1) != 0,
usesFlatAddress);
break;
}
case Gen5ShaderEncoding.Mubuf:
@@ -0,0 +1,106 @@
// 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);
}
@@ -0,0 +1,102 @@
// Copyright (C) 2026 SharpEmu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
using System.Buffers.Binary;
using SharpEmu.ShaderCompiler.Vulkan;
using Xunit;
namespace SharpEmu.Libs.Tests.Agc;
// Structural validation of the GPU detile compute kernel. This cannot run the
// shader without a Vulkan device, but it pins the SPIR-V is well-formed: the
// header is correct, every instruction's word count sums to exactly the module
// length (the classic hand-emit bug), and the compute-specific pieces are
// present (a GLCompute entry point, a LocalSize execution mode, and a runtime
// array for the storage buffers). Full pixel correctness is verified on a GPU.
public sealed class DetileComputeSpirvTests
{
private const uint SpirvMagic = 0x07230203;
private const uint SpirvVersion15 = 0x00010500;
private const ushort OpEntryPoint = 15;
private const ushort OpExecutionMode = 16;
private const ushort OpTypeRuntimeArray = 29;
private const ushort OpFunction = 54;
private const ushort OpFunctionEnd = 56;
private const uint ExecutionModelGLCompute = 5;
private const uint ExecutionModeLocalSize = 17;
[Fact]
public void CreateDetileCompute_EmitsWellFormedComputeModule()
{
var spirv = SpirvFixedShaders.CreateDetileCompute();
Assert.True(spirv.Length % sizeof(uint) == 0, "SPIR-V must be a whole number of words.");
var words = new uint[spirv.Length / sizeof(uint)];
for (var i = 0; i < words.Length; i++)
{
words[i] = BinaryPrimitives.ReadUInt32LittleEndian(spirv.AsSpan(i * sizeof(uint)));
}
Assert.True(words.Length > 5, "Module must have a header plus instructions.");
Assert.Equal(SpirvMagic, words[0]);
Assert.Equal(SpirvVersion15, words[1]);
var bound = words[3];
Assert.True(bound > 1, "Id bound must be set.");
var sawComputeEntry = false;
var sawLocalSize = false;
var sawRuntimeArray = false;
var functionCount = 0;
var functionEndCount = 0;
var offset = 5;
while (offset < words.Length)
{
var word = words[offset];
var wordCount = (int)(word >> 16);
var opcode = (ushort)(word & 0xFFFF);
Assert.True(wordCount >= 1, $"Instruction at {offset} has a zero word count.");
Assert.True(
offset + wordCount <= words.Length,
$"Instruction at {offset} (op {opcode}, wc {wordCount}) overruns the module.");
switch (opcode)
{
case OpEntryPoint when words[offset + 1] == ExecutionModelGLCompute:
sawComputeEntry = true;
break;
case OpExecutionMode
when wordCount >= 6 &&
words[offset + 2] == ExecutionModeLocalSize &&
words[offset + 3] == 8 &&
words[offset + 4] == 8 &&
words[offset + 5] == 1:
sawLocalSize = true;
break;
case OpTypeRuntimeArray:
sawRuntimeArray = true;
break;
case OpFunction:
functionCount++;
break;
case OpFunctionEnd:
functionEndCount++;
break;
}
offset += wordCount;
}
// Word counts must tile the module exactly — a wrong length lands here.
Assert.Equal(words.Length, offset);
Assert.True(sawComputeEntry, "Missing a GLCompute OpEntryPoint.");
Assert.True(sawLocalSize, "Missing an 8x8x1 LocalSize execution mode.");
Assert.True(sawRuntimeArray, "Missing a runtime array (storage buffers).");
Assert.Equal(1, functionCount);
Assert.Equal(1, functionEndCount);
}
}
@@ -83,4 +83,152 @@ public sealed class GnmTilingDetileTests
Assert.Equal((ushort)i, value);
}
}
// GetDetileParams must reproduce TryDetile bit-for-bit: the CPU fallback and
// the GPU compute kernel both consume these params, so a detile driven purely
// by DetileParams (the shared addressing formula the kernel runs) must equal
// the shipped CPU detile for every supported mode/bpp.
[Theory]
[InlineData(27u, 2, 384, 200)] // 64 KiB RB+ R_X (exact-XOR)
[InlineData(27u, 4, 256, 256)] // 64 KiB RB+ R_X (exact-XOR)
[InlineData(9u, 4, 300, 300)] // 64 KiB standard (exact-XOR)
[InlineData(24u, 4, 128, 256)] // 64 KiB RB+ Z_X (exact-XOR)
[InlineData(5u, 4, 200, 120)] // 4 KiB standard (exact-XOR)
[InlineData(8u, 4, 128, 128)] // 64 KiB Z (block-table path)
[InlineData(1u, 4, 64, 64)] // 256 B standard (block-table path)
public void GetDetileParams_ReproducesTryDetile(uint mode, int bpp, int w, int h)
{
var p = GnmTiling.GetDetileParams(mode, bpp, w, h);
Assert.True(p.IsSupported);
// Whole-block tiled buffer (block addressing overshoots the linear extent),
// filled with a deterministic non-trivial pattern.
var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight;
var tiled = new byte[(long)p.BlocksPerRow * blocksHigh * p.BlockBytes];
for (var i = 0; i < tiled.Length; i++)
{
tiled[i] = (byte)((i * 31 + 7) & 0xFF);
}
var expected = new byte[w * h * bpp];
Assert.True(GnmTiling.TryDetile(tiled, expected, mode, w, h, bpp));
var actual = DetileViaParams(tiled, p, w, h, bpp);
Assert.Equal(expected, actual);
}
// The production GnmTiling.DetileWithParams (the active CPU fallback used by
// the Metal path under default-on GPU detile) must equal TryDetile for every
// supported mode/bpp — same DetileParams addressing, no re-derived swizzle.
[Theory]
[InlineData(27u, 2, 384, 200)]
[InlineData(27u, 4, 256, 256)]
[InlineData(9u, 4, 300, 300)]
[InlineData(24u, 4, 128, 256)]
[InlineData(5u, 4, 200, 120)]
[InlineData(8u, 4, 128, 128)]
[InlineData(1u, 4, 64, 64)]
[InlineData(27u, 8, 256, 256)] // 8bpp (GPU: 2 words/element)
[InlineData(27u, 16, 128, 128)] // 16bpp (GPU: 4 words/element)
[InlineData(9u, 8, 128, 96)]
[InlineData(8u, 16, 64, 64)] // block-table, 16bpp
public void DetileWithParams_MatchesTryDetile(uint mode, int bpp, int w, int h)
{
var p = GnmTiling.GetDetileParams(mode, bpp, w, h);
Assert.True(p.IsSupported);
var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight;
var tiled = new byte[(long)p.BlocksPerRow * blocksHigh * p.BlockBytes];
for (var i = 0; i < tiled.Length; i++)
{
tiled[i] = (byte)((i * 31 + 7) & 0xFF);
}
var expected = new byte[w * h * bpp];
Assert.True(GnmTiling.TryDetile(tiled, expected, mode, w, h, bpp));
var actual = new byte[w * h * bpp];
Assert.True(GnmTiling.DetileWithParams(p, tiled, actual));
Assert.Equal(expected, actual);
}
// Array textures are packed as contiguous tiled slices and detiled one slice
// per layer (dispatch-Z on the GPU; a per-layer loop in the CPU fallbacks)
// into a layer-major linear buffer. This pins that packing: each slice must
// deswizzle into its own region and match a per-slice TryDetile, and a
// per-layer-distinct pattern catches any slice cross-talk.
[Theory]
[InlineData(27u, 4, 256, 256, 3)]
[InlineData(9u, 4, 128, 96, 2)]
[InlineData(24u, 4, 64, 128, 4)]
public void DetileWithParams_MultiLayer_MatchesPerSliceTryDetile(uint mode, int bpp, int w, int h, int layers)
{
var p = GnmTiling.GetDetileParams(mode, bpp, w, h);
Assert.True(p.IsSupported);
var blocksHigh = (h + p.BlockHeight - 1) / p.BlockHeight;
var sliceTiledBytes = (int)((long)p.BlocksPerRow * blocksHigh * p.BlockBytes);
var sliceLinearBytes = w * h * bpp;
var tiled = new byte[sliceTiledBytes * layers];
for (var layer = 0; layer < layers; layer++)
{
for (var i = 0; i < sliceTiledBytes; i++)
{
tiled[layer * sliceTiledBytes + i] = (byte)((i * 31 + 7 + layer * 101) & 0xFF);
}
}
// Expected: each slice detiled independently via the shipped CPU detile.
var expected = new byte[sliceLinearBytes * layers];
for (var layer = 0; layer < layers; layer++)
{
Assert.True(GnmTiling.TryDetile(
tiled.AsSpan(layer * sliceTiledBytes, sliceTiledBytes),
expected.AsSpan(layer * sliceLinearBytes, sliceLinearBytes),
mode, w, h, bpp));
}
// Actual: the layer-major loop the Vulkan/Metal CPU fallbacks run.
var actual = new byte[sliceLinearBytes * layers];
for (var layer = 0; layer < layers; layer++)
{
Assert.True(GnmTiling.DetileWithParams(
p,
tiled.AsSpan(layer * sliceTiledBytes, sliceTiledBytes),
actual.AsSpan(layer * sliceLinearBytes, sliceLinearBytes)));
}
Assert.Equal(expected, actual);
}
// Reference detile driven entirely by DetileParams — the single shared
// addressing formula the Vulkan/Metal compute kernel will run per texel.
private static byte[] DetileViaParams(byte[] tiled, DetileParams p, int w, int h, int bpp)
{
var linear = new byte[w * h * bpp];
for (var y = 0; y < h; y++)
{
for (var x = 0; x < w; x++)
{
var blockX = x / p.BlockWidth;
var blockY = y / p.BlockHeight;
var inX = x % p.BlockWidth;
var inY = y % p.BlockHeight;
var inBlockByte = p.Equation == DetileEquation.ExactXor
? p.XByteTerm[x & p.XMask] ^ p.YByteTerm[y & p.YMask]
: p.BlockTable[inY * p.BlockWidth + inX] * p.BytesPerElement;
var srcByte = ((long)blockY * p.BlocksPerRow + blockX) * p.BlockBytes + inBlockByte;
var dstByte = ((long)y * w + x) * bpp;
if (srcByte < 0 || srcByte + bpp > tiled.Length)
{
continue;
}
Array.Copy(tiled, srcByte, linear, dstByte, bpp);
}
}
return linear;
}
}
@@ -0,0 +1,231 @@
// 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,7 +83,9 @@ public sealed class BinkFramePlaybackTests
public uint Height => 1;
public uint FramesPerSecondNumerator => 20;
// 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 FramesPerSecondDenominator => 1;
@@ -11,6 +11,11 @@ public sealed class SystemServiceExportsTests
{
private const ulong MemoryBase = 0x1_0000_0000;
public SystemServiceExportsTests()
{
SystemServiceExports.ResetForTests();
}
[Fact]
public void GetNoticeScreenSkipFlagWritesOneByteAtMemoryBoundary()
{
@@ -26,4 +31,23 @@ 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,4 +36,23 @@ 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));
}
}
@@ -0,0 +1,49 @@
// 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));
}
}
@@ -0,0 +1,179 @@
// 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;
}
}
}
@@ -0,0 +1,197 @@
// 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);
}
@@ -0,0 +1,22 @@
<!--
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>