From 239a644270874cc80bf5ed62967869cadabcae40 Mon Sep 17 00:00:00 2001 From: Codex Date: Tue, 29 Sep 2026 19:13:20 +0300 Subject: [PATCH] Validate native D3D12 HIP interop and isolate Windows mapped-view leak --- CMakeLists.txt | 11 +- docs/windows-rdna4-port.md | 59 ++- .../0001-external-memory-view-ownership.patch | 25 ++ scripts/windows/build-windows-rdna4.ps1 | 6 + scripts/windows/setup-windows-tools.ps1 | 14 +- scripts/windows/test-windows-rdna4.ps1 | 39 +- tools/windows/d3d12_hip_interop_probe.cpp | 346 ++++++++++++++++++ tools/windows/diagnostic.h | 4 + tools/windows/hip_api.h | 3 + tools/windows/hip_stream_lifecycle_probe.cpp | 60 +++ tools/windows/probe_kernel.hip | 12 + 11 files changed, 561 insertions(+), 18 deletions(-) create mode 100644 patches/rocm/0001-external-memory-view-ownership.patch create mode 100644 tools/windows/d3d12_hip_interop_probe.cpp create mode 100644 tools/windows/hip_stream_lifecycle_probe.cpp diff --git a/CMakeLists.txt b/CMakeLists.txt index 87ef6bc..721d595 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -12,7 +12,7 @@ endif() if(NOT EXISTS "${D4R_HIP_ROOT}/include/hip/hip_runtime_api.h") message(FATAL_ERROR "Set D4R_HIP_ROOT to a Windows HIP SDK (headers and bin/clang++.exe required)") endif() -find_program(D4R_HIP_CLANG NAMES clang++.exe PATHS "${D4R_HIP_ROOT}/bin" NO_DEFAULT_PATH REQUIRED) +find_program(D4R_HIP_CLANG NAMES clang++.exe PATHS "${D4R_HIP_ROOT}/bin" "${D4R_HIP_ROOT}/lib/llvm/bin" NO_DEFAULT_PATH REQUIRED) set(probe_object "${CMAKE_CURRENT_BINARY_DIR}/probe_gfx1201.hsaco") add_custom_command(OUTPUT "${probe_object}" COMMAND "${D4R_HIP_CLANG}" -x hip --offload-arch=gfx1201 @@ -21,7 +21,7 @@ add_custom_command(OUTPUT "${probe_object}" -o "${probe_object}" DEPENDS tools/windows/probe_kernel.hip VERBATIM) add_custom_target(d4r_probe_code_object ALL DEPENDS "${probe_object}") -foreach(probe hip_gfx1201_probe cuda_driver_probe) +foreach(probe hip_gfx1201_probe cuda_driver_probe d3d12_hip_interop_probe hip_stream_lifecycle_probe) add_executable(d4r_${probe} tools/windows/${probe}.cpp) target_compile_definitions(d4r_${probe} PRIVATE WIN32_LEAN_AND_MEAN NOMINMAX __HIP_PLATFORM_AMD__) target_include_directories(d4r_${probe} PRIVATE "${D4R_HIP_ROOT}/include") @@ -37,6 +37,7 @@ foreach(probe hip_gfx1201_probe cuda_driver_probe) endif() install(TARGETS d4r_${probe} RUNTIME DESTINATION bin) endforeach() +target_link_libraries(d4r_d3d12_hip_interop_probe PRIVATE d3d12 dxgi dxguid) add_dependencies(d4r_hip_gfx1201_probe d4r_probe_code_object) install(FILES "${probe_object}" DESTINATION bin) install(FILES scripts/windows/test-windows-rdna4.ps1 DESTINATION .) @@ -45,6 +46,12 @@ include(CTest) add_test(NAME hip_gfx1201 COMMAND d4r_hip_gfx1201_probe --hip-root "${D4R_HIP_ROOT}" --module "${probe_object}") set_tests_properties(hip_gfx1201 PROPERTIES TIMEOUT 90 LABELS "hardware;gfx1201") +add_test(NAME d3d12_hip_interop COMMAND d4r_d3d12_hip_interop_probe + --hip-root "${D4R_HIP_ROOT}" --module "${probe_object}" --iterations 32) +set_tests_properties(d3d12_hip_interop PROPERTIES TIMEOUT 120 LABELS "hardware;interop") +add_test(NAME d3d12_hip_map_lifetime COMMAND d4r_d3d12_hip_interop_probe + --hip-root "${D4R_HIP_ROOT}" --interop-mode map --iterations 64) +set_tests_properties(d3d12_hip_map_lifetime PROPERTIES TIMEOUT 120 LABELS "hardware;interop;lifetime") set(D4R_ZLUDA_ROOT "" CACHE PATH "Local Windows ZLUDA directory (nvcuda.dll)") if(EXISTS "${D4R_ZLUDA_ROOT}/nvcuda.dll") foreach(context primary created) diff --git a/docs/windows-rdna4-port.md b/docs/windows-rdna4-port.md index 35dd9f1..3a47c5d 100644 --- a/docs/windows-rdna4-port.md +++ b/docs/windows-rdna4-port.md @@ -6,8 +6,8 @@ This is a work log for an implementation in progress, not a claim of DLSS suppor ## Milestone and gates -Current milestone: M1/M2 passed on the physical RX 9070 XT; next: independent -Windows D3D12/HIP interop and validated gfx12 WMMA, before NGX integration. +Current milestone: M1/M2 and M4 passed on the physical RX 9070 XT with the +isolated TheRock profile. Next: NGX initialization and validated gfx12 WMMA. Do not integrate NGX until the integer PTX workload is stable on the real GPU. | Gate | Status | @@ -16,7 +16,7 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. | M1: native HIP gfx1201 allocation, kernel, CPU verification | PASS, 32 iterations + guard verification | | M2: Windows ZLUDA integer PTX, primary and created contexts | PASS, 32 iterations each + guard verification | | M3: Windows NGX initialization and simple DLSS path | Pending | -| M4: D3D12 / HIP external memory and fence round trip | Pending | +| M4: D3D12 / HIP external memory and fence round trip | PASS with TheRock; stable 7.2 has mapped-view leak | | M5: independently validated gfx12 WMMA backend | Pending | | M6: K layers, full transformer and image validation | Pending | | M7: M FP16-equivalent baseline and full transformer | Pending | @@ -36,6 +36,10 @@ Do not integrate NGX until the integer PTX workload is stable on the real GPU. LLVM commit `590b9320a5be90e40268759c6203c01fde121e68`. * Keep stable SDK and TheRock installations separate. Select one directory per build and test run; never mix headers, import libraries, runtime or COMGR. +* Tested isolated TheRock core: `10.2.0a20260929`, HIP `7.17.26386`, commit + `0ad5f73253`; wheel SHA256 + `9623b97ca511eaa176905075a504393eca6fde594222af10b55e4fb1edfe5a25`. + No machine-wide runtime replacement is performed. * CMake 3.31.6, Ninja 1.11.1 (Python package 1.11.1.4). * Windows host toolchain: LLVM-MinGW `20260922`, UCRT x64, SHA256 `e3ad77d117a4bea19a7a3b333341824d79a5a371004a10e25b8504e7b3047666`. @@ -125,6 +129,25 @@ NT HANDLE ownership must be explicit; Win32 import is not fd ownership transfer. Test repeated create/destroy cycles as upstream noted a mapped-buffer leak on Linux ROCm. CPU verification copies are allowed in tests, never in the fast path. +Implemented independent test: DXGI LUID matches HIP, an R32_UINT texture is +copied to a shared DEFAULT-heap linear buffer in VRAM, HIP maps that D3D12 +resource, checks every input element and XORs it, D3D12 copies it back to the +texture and finally to a readback buffer solely for test verification. Shared +D3D12 fence values sequence both queues within the same iteration. There is no +CPU data copy between the APIs; native tiled texture import is not yet tested. +The future shim needs GPU texture/linear conversion or validated direct arrays. +The caller closes its original NT HANDLEs after HIP releases imports. + +Stable HIP SDK `7.2.60201` fails mapped-buffer lifetime testing. `resource` and +`import` modes are stable, but `map` leaks one HANDLE and 262144 bytes per cycle. +The runtime's extra `view->retain()` is the cause; AMD removed it in +[CLR commit 529f6b1](https://github.com/ROCm/clr/commit/529f6b1641de436dafbb095d5438b0fbf773765d). +The source backport is retained at `patches/rocm/0001-external-memory-view-ownership.patch`. +No pointer/refcount or DLL binary hacks are used. Current TheRock passes both +the 64-cycle map lifetime test and the 32-cycle synchronized texture round trip. +The runner automatically isolates resource/import/stream lifetimes on failure. +Do not declare stable SDK 7.2 acceptable for the leak-free fast path. + ## Authoritative references * [AMD HIP SDK Windows support matrix](https://rocm.docs.amd.com/projects/install-on-windows/en/latest/reference/system-requirements.html) @@ -151,6 +174,18 @@ powershell -NoProfile -File scripts/windows/build-windows-rdna4.ps1 powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 ``` +For the currently validated external-memory runtime, use the separate profile: + +```powershell +powershell -NoProfile -File scripts/windows/setup-windows-tools.ps1 -RuntimeProfile therock +powershell -NoProfile -File scripts/windows/build-windows-rdna4.ps1 -RuntimeProfile therock +powershell -NoProfile -File scripts/windows/test-windows-rdna4.ps1 -RuntimeProfile therock +``` + +This builds into `build/windows-rdna4-therock` and installs into +`dist/windows-rdna4-therock`. The stable profile stays available for reproducing +the runtime bug. It is expected to return failure at the lifetime gate. + The build accepts `-HipRoot`, `-ZludaRoot`, `-ToolchainRoot`, `-BuildDirectory` and `-InstallDirectory`. MSVC users can invoke CMake from a Developer Shell without the portable MinGW compiler. Equivalent direct build: @@ -165,11 +200,13 @@ cmake --install build/windows-rdna4 --prefix dist/windows-rdna4-diagnostics The one-command diagnostic runner is also installed at `dist/windows-rdna4-diagnostics/test-windows-rdna4.ps1`. It requires the selected HIP SDK and ZLUDA directory; the developer setup auto-discovers local defaults. -Success: exit 0, all three tests passed in `summary.json`, `PASS HIP` and -`PASS CUDA` in stdout. Failure: send the **single printed ZIP path**; it contains +Success: exit 0, all five tests passed in `summary.json`, `PASS HIP`, `PASS CUDA`, +`PASS INTEROP_LIFETIME` and `PASS INTEROP` in stdout. Failure: send the **single printed ZIP path**; it contains stdout/stderr, timeout/exception exit code, DLL paths/versions/hashes, driver/OS inventory and a minidump for an unhandled exception. Upstream trace is attempted -automatically on CUDA failure when `zluda.exe` is available. +automatically on CUDA failure when `trace/nvcuda.dll` is available. The trace +DLL forwards to the real DLL through `ZLUDA_CUDA_LIB`; explicit loading no +longer bypasses the trace by accidentally passing the real DLL to the launcher. The HIP module compiler uses the selected SDK, fixed `--offload-arch=gfx1201`, wave32 and no host/device library dependency. Host code uses the installed HIP @@ -194,6 +231,14 @@ all context/module/memory teardown calls successful. ZLUDA binary SHA256 `51dd32dc116a6c14c7a4bf5acf620bba0546ca7ec5f8a6e2148dd68f1714b165`. Logs: local `test-results/m1-m2.zip`; ignored by git. +2026-09-29 M4: TheRock profile passed all HIP/PTX checks plus D3D12/HIP round +trip (32 iterations, 16384 elements each), exact integer output and GPU fence +synchronization. Map lifetime: 64 import/map/free/destroy cycles, HANDLE count +267 after both warmup and the last cycle, free VRAM 16926011392 bytes throughout +the measured steady state. Full interop cycle HANDLE count also stays constant +after warmup; stream teardown lowers it. Recreating streams did not resolve the +stable runtime mapped-view bug; the fix is the runtime ownership correction. + Local proprietary inputs supplied by the user: `nvngx_dlss.dll` 310.9.1.0, `_nvngx.dll` 32.0.16.1714. Neither is packaged or tracked. @@ -206,4 +251,4 @@ launching the MinGW child compiler, so branch/commit/build require scoped tool approval in this managed environment. HIP_PATH ends in a backslash; build script normalizes it before PowerShell 5 argument quoting. Test runner uses .NET Process to retain reliable exit codes and concurrently drain both diagnostic streams. -No K/M/DLSS/OptiScaler support is claimed by this first milestone. +No K/M/DLSS/OptiScaler support is claimed by these diagnostic milestones. diff --git a/patches/rocm/0001-external-memory-view-ownership.patch b/patches/rocm/0001-external-memory-view-ownership.patch new file mode 100644 index 0000000..cb438fa --- /dev/null +++ b/patches/rocm/0001-external-memory-view-ownership.patch @@ -0,0 +1,25 @@ +From: d4r Windows port +Subject: [PATCH] Backport AMD external-memory view ownership fix to ROCm 7.2 + +Upstream fix: ROCm/clr 529f6b1641de436dafbb095d5438b0fbf773765d +(ROCm/rocm-systems #3528). Already present in current upstream. +The constructor owns reference 1. MemObjMap does not retain the view; +hipFree removes the entry and releases reference 1. The extra retain +leaks the view and its parent imported allocation. + +Hardware regression: d4r_d3d12_hip_interop_probe --interop-mode map +--iterations 64. Windows SDK 7.2.60201 leaks one HANDLE / 256 KiB each +cycle; tested TheRock 10.2.0a20260929 (HIP 7.17.26386) stays bounded. +This patch is for rebuilding old runtime sources, never binary patching. + +diff --git a/hipamd/src/hip_memory.cpp b/hipamd/src/hip_memory.cpp +--- a/hipamd/src/hip_memory.cpp ++++ b/hipamd/src/hip_memory.cpp +@@ -179,7 +179,6 @@ hipError_t hipExternalMemoryGetMappedBuffer(void** devPtr, hipExternalMemory_t extMem, + // Map the device memory to the user pointer + *devPtr = reinterpret_cast(devMem->virtualAddress()); + amd::MemObjMap::AddMemObj(*devPtr, view); +- view->retain(); + + HIP_RETURN(hipSuccess); + } diff --git a/scripts/windows/build-windows-rdna4.ps1 b/scripts/windows/build-windows-rdna4.ps1 index 3f7ea05..10fa2d2 100644 --- a/scripts/windows/build-windows-rdna4.ps1 +++ b/scripts/windows/build-windows-rdna4.ps1 @@ -1,6 +1,7 @@ [CmdletBinding()] param( [string]$HipRoot = $env:HIP_PATH, + [ValidateSet('stable', 'therock')][string]$RuntimeProfile = 'stable', [string]$ZludaRoot, [string]$ToolchainRoot, [string]$BuildDirectory, @@ -8,6 +9,11 @@ param( ) $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) +if ($RuntimeProfile -eq 'therock') { + if (!$PSBoundParameters.ContainsKey('HipRoot')) { $HipRoot = Join-Path $repo '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } + if (!$BuildDirectory) { $BuildDirectory = Join-Path $repo 'build/windows-rdna4-therock' } + if (!$InstallDirectory) { $InstallDirectory = Join-Path $repo 'dist/windows-rdna4-therock' } +} if (!$HipRoot) { $HipRoot = 'C:\Program Files\AMD\ROCm\7.2' } if (!$BuildDirectory) { $BuildDirectory = Join-Path $repo 'build/windows-rdna4' } if (!$InstallDirectory) { $InstallDirectory = Join-Path $repo 'dist/windows-rdna4-diagnostics' } diff --git a/scripts/windows/setup-windows-tools.ps1 b/scripts/windows/setup-windows-tools.ps1 index b8a1c04..6aab3ae 100644 --- a/scripts/windows/setup-windows-tools.ps1 +++ b/scripts/windows/setup-windows-tools.ps1 @@ -1,5 +1,5 @@ [CmdletBinding()] -param() +param([ValidateSet('stable', 'therock')][string]$RuntimeProfile = 'stable') $ErrorActionPreference = 'Stop' $repo = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) $localTools = Join-Path $repo '.tools' @@ -14,7 +14,10 @@ function Get-VerifiedArchive([string]$Url, [string]$Name, [string]$Hash, [string if ((Get-FileHash -LiteralPath $archive -Algorithm SHA256).Hash -ne $Hash) { throw "SHA256 mismatch: $archive (remove this download and retry)" } - Expand-Archive -LiteralPath $archive -DestinationPath $Destination -Force + if ([IO.Path]::GetExtension($archive) -eq '.whl') { + python -c 'import sys,zipfile; zipfile.ZipFile(sys.argv[1]).extractall(sys.argv[2])' $archive $Destination + if ($LASTEXITCODE) { throw "Wheel extraction exit=$LASTEXITCODE" } + } else { Expand-Archive -LiteralPath $archive -DestinationPath $Destination -Force } if (!(Test-Path -LiteralPath $Marker)) { throw "Archive layout changed: $Marker" } } Get-VerifiedArchive 'https://github.com/mstorsjo/llvm-mingw/releases/download/20260922/llvm-mingw-20260922-ucrt-x86_64.zip' ` @@ -23,6 +26,13 @@ Get-VerifiedArchive 'https://github.com/mstorsjo/llvm-mingw/releases/download/20 Get-VerifiedArchive 'https://github.com/vosen/ZLUDA/releases/download/v7-preview.11/zluda-windows-ee2f25a.zip' ` 'zluda.zip' '7788c1ed43385e62f0cc5f4d1aead24c92b671c02f1d685f63182540dea16d74' (Join-Path $localTools 'zluda') ` (Join-Path $localTools 'zluda/zluda/nvcuda.dll') +if ($RuntimeProfile -eq 'therock') { + Get-VerifiedArchive 'https://nightly.repo.amd.com/rocm/core/whl-next/rocm-sdk-core/rocm_sdk_core-10.2.0a20260929-py3-none-win_amd64.whl' ` + 'rocm_sdk_core-10.2.0a20260929-py3-none-win_amd64.whl' ` + '9623b97ca511eaa176905075a504393eca6fde594222af10b55e4fb1edfe5a25' ` + (Join-Path $localTools 'therock-10.2.0a20260929') ` + (Join-Path $localTools 'therock-10.2.0a20260929/_rocm_sdk_core/bin/amdhip64_7.dll') +} if (!(Test-Path (Join-Path $localTools 'python/cmake/data/bin/cmake.exe')) -or !(Test-Path (Join-Path $localTools 'python/bin/ninja.exe'))) { python -m pip install --target (Join-Path $localTools 'python') cmake==3.31.6 ninja==1.11.1.4 diff --git a/scripts/windows/test-windows-rdna4.ps1 b/scripts/windows/test-windows-rdna4.ps1 index b94caf4..8e04c17 100644 --- a/scripts/windows/test-windows-rdna4.ps1 +++ b/scripts/windows/test-windows-rdna4.ps1 @@ -1,6 +1,7 @@ [CmdletBinding()] param( [string]$HipRoot = $env:HIP_PATH, + [ValidateSet('stable', 'therock')][string]$RuntimeProfile = 'stable', [string]$ZludaRoot, [string]$PackageRoot, [string]$OutputDirectory, @@ -8,6 +9,15 @@ param( [int]$TimeoutSeconds = 180 ) $ErrorActionPreference = 'Stop' +if ($RuntimeProfile -eq 'therock') { + $repoRoot = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../..')) + # Works both from scripts/windows and the installed dist/ directory. + if (!$PSBoundParameters.ContainsKey('HipRoot')) { $HipRoot = Join-Path $repoRoot '.tools/therock-10.2.0a20260929/_rocm_sdk_core' } + if (!$PackageRoot) { + if (Test-Path (Join-Path $PSScriptRoot 'bin/d4r_hip_gfx1201_probe.exe')) { $PackageRoot = $PSScriptRoot } + else { $PackageRoot = Join-Path $repoRoot 'dist/windows-rdna4-therock' } + } +} if (!$PackageRoot) { if (Test-Path (Join-Path $PSScriptRoot 'bin/d4r_hip_gfx1201_probe.exe')) { $PackageRoot = $PSScriptRoot } else { $PackageRoot = [IO.Path]::GetFullPath((Join-Path $PSScriptRoot '../../dist/windows-rdna4-diagnostics')) } @@ -30,7 +40,7 @@ $originalDiag = $env:D4R_DIAG_DIR $originalLog = $env:ZLUDA_LOG_DIR $originalCuda = $env:ZLUDA_CUDA_LIB $exitStatus = 1 -$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()} +$summary = [ordered]@{utc=[DateTime]::UtcNow.ToString('o'); profile=$RuntimeProfile; hipRoot=$HipRoot; zludaRoot=$ZludaRoot; tests=@()} function Quote-Argument([string]$Value) { # All generated paths are absolute file/directory paths, with no trailing backslash. if ($Value.Contains('"')) { throw 'Double quotes are not allowed in arguments' } @@ -76,7 +86,9 @@ try { $env:D4R_DIAG_DIR = $OutputDirectory $env:ZLUDA_LOG_DIR = Join-Path $OutputDirectory 'zluda-trace' $env:ZLUDA_CUDA_LIB = Join-Path $ZludaRoot 'nvcuda.dll' - $files = @((Join-Path $HipRoot 'bin/amdhip64_7.dll'), (Join-Path $HipRoot 'bin/amd_comgr_3.dll'), + $comgr = Join-Path $HipRoot 'bin/amd_comgr_3.dll' + if (!(Test-Path -LiteralPath $comgr)) { $comgr = Join-Path $HipRoot 'bin/amd_comgr.dll' } + $files = @((Join-Path $HipRoot 'bin/amdhip64_7.dll'), $comgr, (Join-Path $ZludaRoot 'nvcuda.dll'), (Join-Path $PackageRoot 'bin/probe_gfx1201.hsaco')) $summary.files = @($files | ForEach-Object { if (!(Test-Path -LiteralPath $_)) { throw "Required file missing: $_" } @@ -100,14 +112,27 @@ try { '--context', $context, '--iterations', "$Iterations") $cudaOk = $cudaOk -and $ok # On failure collect the upstream trace automatically if available. - if (!$ok -and (Test-Path (Join-Path $ZludaRoot 'zluda.exe'))) { - Invoke-Probe "cuda-$context-trace" (Join-Path $ZludaRoot 'zluda.exe') @( - '--zluda-trace', '--', (Join-Path $bin 'd4r_cuda_driver_probe.exe'), - '--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'nvcuda.dll'), + if (!$ok -and (Test-Path (Join-Path $ZludaRoot 'trace/nvcuda.dll'))) { + Invoke-Probe "cuda-$context-trace" (Join-Path $bin 'd4r_cuda_driver_probe.exe') @( + '--hip-root', $HipRoot, '--cuda-dll', (Join-Path $ZludaRoot 'trace/nvcuda.dll'), '--context', $context, '--iterations', '1') | Out-Null } } - if ($cudaOk) { $exitStatus = 0 } + if ($cudaOk) { + $mapOk = Invoke-Probe 'interop-map-lifetime' (Join-Path $bin 'd4r_d3d12_hip_interop_probe.exe') @( + '--hip-root', $HipRoot, '--interop-mode', 'map', '--iterations', '64') + $interopOk = Invoke-Probe 'interop-roundtrip' (Join-Path $bin 'd4r_d3d12_hip_interop_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $bin 'probe_gfx1201.hsaco'), '--iterations', "$Iterations") + if (!$mapOk -or !$interopOk) { + foreach ($mode in @('resource', 'import')) { + Invoke-Probe "interop-$mode-lifetime" (Join-Path $bin 'd4r_d3d12_hip_interop_probe.exe') @( + '--hip-root', $HipRoot, '--interop-mode', $mode, '--iterations', '32') | Out-Null + } + Invoke-Probe 'hip-stream-lifetime' (Join-Path $bin 'd4r_hip_stream_lifecycle_probe.exe') @( + '--hip-root', $HipRoot, '--module', (Join-Path $bin 'probe_gfx1201.hsaco'), '--iterations', '32') | Out-Null + } + if ($mapOk -and $interopOk) { $exitStatus = 0 } + } } } catch { $summary.error = $_.Exception.Message diff --git a/tools/windows/d3d12_hip_interop_probe.cpp b/tools/windows/d3d12_hip_interop_probe.cpp new file mode 100644 index 0000000..14d289d --- /dev/null +++ b/tools/windows/d3d12_hip_interop_probe.cpp @@ -0,0 +1,346 @@ +#if defined(__MINGW32__) +#define WIDL_EXPLICIT_AGGREGATE_RETURNS +#endif +#include "hip_api.h" +#include +#include +#include + +using Microsoft::WRL::ComPtr; +using namespace d4r::diag; + +static void handle_checkpoint(const char* stage) +{ + DWORD count = 0; + if (!GetProcessHandleCount(GetCurrentProcess(), &count)) + throw std::runtime_error("GetProcessHandleCount failed"); + std::printf("HANDLE_STAGE %s count=%lu\n", stage, count); +} + +static void dx(HRESULT hr, const char* call) +{ + std::printf("D3D12 %s -> 0x%08lx\n", call, static_cast(hr)); + if (FAILED(hr)) throw std::runtime_error(std::string(call) + " HRESULT=" + std::to_string(hr)); +} +struct NativeHandle { + HANDLE handle = nullptr; + ~NativeHandle() { if (handle) CloseHandle(handle); } +}; +struct ExternalApi { + HipApi& hip; +#define EXT(name) decltype(&::name) name = hip.library.symbol(#name) + EXT(hipImportExternalMemory); + EXT(hipExternalMemoryGetMappedBuffer); + EXT(hipDestroyExternalMemory); + EXT(hipImportExternalSemaphore); + EXT(hipDestroyExternalSemaphore); + EXT(hipWaitExternalSemaphoresAsync); + EXT(hipSignalExternalSemaphoresAsync); +#undef EXT +}; +struct Imported { + ExternalApi& api; + hipExternalMemory_t memory = nullptr; + hipExternalSemaphore_t semaphore = nullptr; + hipStream_t stream = nullptr; + bool owns_stream = true; + void* mapped = nullptr; + void* errors = nullptr; + ~Imported() { + // No resource may be released while the HIP stream is still using it. + if (stream) (void)api.hip.hipStreamSynchronize(stream); + if (stream && owns_stream) (void)api.hip.hipStreamDestroy(stream); + if (errors) (void)api.hip.hipFree(errors); + if (mapped) (void)api.hip.hipFree(mapped); + if (memory) (void)api.hipDestroyExternalMemory(memory); + if (semaphore) (void)api.hipDestroyExternalSemaphore(semaphore); + } +}; +static D3D12_RESOURCE_DESC buffer_desc(uint64_t bytes) +{ + D3D12_RESOURCE_DESC desc{}; + desc.Dimension = D3D12_RESOURCE_DIMENSION_BUFFER; + desc.Width = bytes; desc.Height = 1; desc.DepthOrArraySize = 1; + desc.MipLevels = 1; desc.SampleDesc.Count = 1; + desc.Layout = D3D12_TEXTURE_LAYOUT_ROW_MAJOR; + return desc; +} +static ComPtr make_buffer(ID3D12Device* device, uint64_t bytes, + D3D12_HEAP_TYPE type, D3D12_HEAP_FLAGS flags, D3D12_RESOURCE_STATES state) +{ + D3D12_HEAP_PROPERTIES properties{}; + properties.Type = type; + properties.CreationNodeMask = properties.VisibleNodeMask = 1; + auto desc = buffer_desc(bytes); + ComPtr resource; + dx(device->CreateCommittedResource(&properties, flags, &desc, state, nullptr, + IID_PPV_ARGS(resource.GetAddressOf())), "CreateCommittedResource(buffer)"); + return resource; +} +static void transition(ID3D12GraphicsCommandList* list, ID3D12Resource* resource, + D3D12_RESOURCE_STATES before, D3D12_RESOURCE_STATES after) +{ + D3D12_RESOURCE_BARRIER barrier{}; + barrier.Type = D3D12_RESOURCE_BARRIER_TYPE_TRANSITION; + barrier.Transition.pResource = resource; + barrier.Transition.Subresource = D3D12_RESOURCE_BARRIER_ALL_SUBRESOURCES; + barrier.Transition.StateBefore = before; barrier.Transition.StateAfter = after; + list->ResourceBarrier(1, &barrier); +} + +static void import_lifetime(ID3D12Device* device, ExternalApi& api, const Args& args) +{ + if (args.interop_mode != "resource" && args.interop_mode != "import" && args.interop_mode != "map") + throw std::runtime_error("--interop-mode must be roundtrip, resource, import or map"); + DWORD baseline = 0, handles = 0; + for (unsigned i = 0; i < args.iterations; ++i) { + { + auto resource = make_buffer(device, 65536, D3D12_HEAP_TYPE_DEFAULT, + D3D12_HEAP_FLAG_SHARED, D3D12_RESOURCE_STATE_COMMON); + NativeHandle handle; + dx(device->CreateSharedHandle(resource.Get(), nullptr, GENERIC_ALL, nullptr, + &handle.handle), "CreateSharedHandle(lifecycle)"); + Imported imported{api}; + if (args.interop_mode != "resource") { + hipExternalMemoryHandleDesc desc{}; + desc.type = hipExternalMemoryHandleTypeD3D12Resource; + desc.handle.win32.handle = handle.handle; + desc.size = 65536; + desc.flags = hipExternalMemoryDedicated; + api.hip.check(api.hipImportExternalMemory(&imported.memory, &desc), "hipImportExternalMemory(lifecycle)"); + if (args.interop_mode == "map") { + hipExternalMemoryBufferDesc map{}; map.size = 65536; + api.hip.check(api.hipExternalMemoryGetMappedBuffer(&imported.mapped, imported.memory, &map), "hipExternalMemoryGetMappedBuffer(lifecycle)"); + api.hip.check(api.hip.hipFree(imported.mapped), "hipFree(lifecycle)"); imported.mapped = nullptr; + } + api.hip.check(api.hipDestroyExternalMemory(imported.memory), "hipDestroyExternalMemory(lifecycle)"); imported.memory = nullptr; + } + } + if (!GetProcessHandleCount(GetCurrentProcess(), &handles)) throw std::runtime_error("GetProcessHandleCount failed"); + size_t free_bytes = 0, total_bytes = 0; + api.hip.check(api.hip.hipMemGetInfo(&free_bytes, &total_bytes), "hipMemGetInfo"); + if (i == 1 || (args.iterations == 1 && i == 0)) baseline = handles; + std::printf("IMPORT_CLEANUP mode=%s iteration=%u handles=%lu baseline=%lu free_vram=%zu\n", + args.interop_mode.c_str(), i, handles, baseline, free_bytes); + } + if (handles > baseline + 4) throw std::runtime_error("External memory lifecycle handle growth"); + std::printf("PASS INTEROP_LIFETIME mode=%s iterations=%u\n", args.interop_mode.c_str(), args.iterations); +} + +int main(int argc, char** argv) +{ + start(); + try { + Args args(argc, argv); + if (args.module.empty() && args.interop_mode == "roundtrip") throw std::runtime_error("--module required"); + if (args.stream_lifetime != "persistent" && args.stream_lifetime != "cycle") + throw std::runtime_error("--stream-lifetime must be persistent or cycle"); + HipApi hip(args.hip_root); + hipDeviceProp_t props{}; + hip.select_gfx1201(args.device, props); + ExternalApi external{hip}; + ComPtr factory; + dx(CreateDXGIFactory1(IID_PPV_ARGS(factory.GetAddressOf())), "CreateDXGIFactory1"); + ComPtr adapter; + bool matched = false; + for (UINT i = 0; ; ++i) { + const HRESULT enumerated = factory->EnumAdapters1(i, adapter.ReleaseAndGetAddressOf()); + if (enumerated == DXGI_ERROR_NOT_FOUND) break; + dx(enumerated, "EnumAdapters1"); + DXGI_ADAPTER_DESC1 desc{}; + dx(adapter->GetDesc1(&desc), "GetDesc1"); + std::printf("DXGI adapter=%s LUID=%08lx:%08lx\n", utf8(desc.Description).c_str(), + static_cast(desc.AdapterLuid.HighPart), desc.AdapterLuid.LowPart); + if (std::memcmp(&desc.AdapterLuid, props.luid, sizeof(LUID)) == 0) { matched = true; break; } + } + if (!matched) throw std::runtime_error("HIP/DXGI LUID mismatch: cross-adapter import prohibited"); + ComPtr device; + dx(D3D12CreateDevice(adapter.Get(), D3D_FEATURE_LEVEL_12_0, IID_PPV_ARGS(device.GetAddressOf())), "D3D12CreateDevice"); + if (args.interop_mode != "roundtrip") { import_lifetime(device.Get(), external, args); return 0; } + ComPtr queue; + D3D12_COMMAND_QUEUE_DESC qd{}; + qd.Type = D3D12_COMMAND_LIST_TYPE_DIRECT; + dx(device->CreateCommandQueue(&qd, IID_PPV_ARGS(queue.GetAddressOf())), "CreateCommandQueue"); + ComPtr shared_fence, complete; + dx(device->CreateFence(0, D3D12_FENCE_FLAG_SHARED, IID_PPV_ARGS(shared_fence.GetAddressOf())), "CreateFence(shared)"); + dx(device->CreateFence(0, D3D12_FENCE_FLAG_NONE, IID_PPV_ARGS(complete.GetAddressOf())), "CreateFence(completion)"); + NativeHandle fence_handle, event; + event.handle = CreateEventW(nullptr, FALSE, FALSE, nullptr); + if (!event.handle) throw std::runtime_error("CreateEvent failed"); + dx(device->CreateSharedHandle(shared_fence.Get(), nullptr, GENERIC_ALL, nullptr, + &fence_handle.handle), "CreateSharedHandle(fence)"); + + // R32_UINT texture -> linear shared buffer -> HIP -> texture -> verification readback. + constexpr UINT width = 256, height = 64, count = width * height; + const uint64_t bytes = static_cast(count) * sizeof(uint32_t); + auto upload = make_buffer(device.Get(), bytes, D3D12_HEAP_TYPE_UPLOAD, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_GENERIC_READ); + auto readback = make_buffer(device.Get(), bytes, D3D12_HEAP_TYPE_READBACK, D3D12_HEAP_FLAG_NONE, D3D12_RESOURCE_STATE_COPY_DEST); + D3D12_RESOURCE_DESC td{}; + td.Dimension = D3D12_RESOURCE_DIMENSION_TEXTURE2D; + td.Width = width; td.Height = height; td.DepthOrArraySize = 1; td.MipLevels = 1; + td.SampleDesc.Count = 1; td.Format = DXGI_FORMAT_R32_UINT; + D3D12_HEAP_PROPERTIES hp{}; hp.Type = D3D12_HEAP_TYPE_DEFAULT; + hp.CreationNodeMask = hp.VisibleNodeMask = 1; + ComPtr texture; + dx(device->CreateCommittedResource(&hp, D3D12_HEAP_FLAG_NONE, &td, D3D12_RESOURCE_STATE_COPY_DEST, + nullptr, IID_PPV_ARGS(texture.GetAddressOf())), "CreateCommittedResource(texture)"); + D3D12_PLACED_SUBRESOURCE_FOOTPRINT footprint{}; + UINT rows; UINT64 row_bytes, total_bytes; + device->GetCopyableFootprints(&td, 0, 1, 0, &footprint, &rows, &row_bytes, &total_bytes); + if (total_bytes != bytes || footprint.Footprint.RowPitch != width * 4) + throw std::runtime_error("Unexpected footprint; test requires no padding"); + auto texture_copy = D3D12_TEXTURE_COPY_LOCATION{}; + texture_copy.pResource = texture.Get(); + texture_copy.Type = D3D12_TEXTURE_COPY_TYPE_SUBRESOURCE_INDEX; + auto linear_copy = D3D12_TEXTURE_COPY_LOCATION{}; + linear_copy.Type = D3D12_TEXTURE_COPY_TYPE_PLACED_FOOTPRINT; + linear_copy.PlacedFootprint = footprint; + hipModule_t module = nullptr; + hip.check(hip.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad"); + struct ModuleCleanup { HipApi& hip; hipModule_t m; ~ModuleCleanup() { (void)hip.hipModuleUnload(m); } } mc{hip, module}; + hipFunction_t kernel = nullptr; + hip.check(hip.hipModuleGetFunction(&kernel, module, "d4r_probe_interop"), "hipModuleGetFunction"); + struct StreamCleanup { + HipApi& hip; + hipStream_t stream = nullptr; + ~StreamCleanup() { if (stream) { + (void)hip.hipStreamSynchronize(stream); + (void)hip.hipStreamDestroy(stream); + } } + } persistent{hip}; + if (args.stream_lifetime == "persistent") + hip.check(hip.hipStreamCreateWithFlags(&persistent.stream, hipStreamNonBlocking), "hipStreamCreateWithFlags(persistent)"); + loaded_modules(); + DWORD handles_before = 0; GetProcessHandleCount(GetCurrentProcess(), &handles_before); + DWORD warmed_handles = 0; + DWORD final_cycle_handles = 0; + for (unsigned iteration = 0; iteration < args.iterations; ++iteration) { + const uint32_t seed = 0x12345678u + iteration * 31u; + auto shared = make_buffer(device.Get(), bytes, D3D12_HEAP_TYPE_DEFAULT, D3D12_HEAP_FLAG_SHARED, + D3D12_RESOURCE_STATE_COPY_DEST); + const auto shared_desc = buffer_desc(bytes); + D3D12_RESOURCE_ALLOCATION_INFO allocation_info{}; +#if defined(__MINGW32__) + device->GetResourceAllocationInfo(&allocation_info, 0, 1, &shared_desc); +#else + allocation_info = device->GetResourceAllocationInfo(0, 1, &shared_desc); +#endif + NativeHandle memory_handle; + dx(device->CreateSharedHandle(shared.Get(), nullptr, GENERIC_ALL, nullptr, + &memory_handle.handle), "CreateSharedHandle(resource)"); + Imported imported{external}; + handle_checkpoint("created_resource"); + hipExternalMemoryHandleDesc memory_desc{}; + memory_desc.type = hipExternalMemoryHandleTypeD3D12Resource; + memory_desc.handle.win32.handle = memory_handle.handle; + memory_desc.size = allocation_info.SizeInBytes; + memory_desc.flags = hipExternalMemoryDedicated; + std::printf("STAGE hipImportExternalMemory(D3D12Resource) size=%llu\n", memory_desc.size); + hip.check(external.hipImportExternalMemory(&imported.memory, &memory_desc), "hipImportExternalMemory(D3D12Resource)"); + handle_checkpoint("imported_memory"); + hipExternalMemoryBufferDesc map_desc{}; map_desc.size = bytes; + hip.check(external.hipExternalMemoryGetMappedBuffer(&imported.mapped, imported.memory, &map_desc), "hipExternalMemoryGetMappedBuffer"); + handle_checkpoint("mapped_memory"); + hipExternalSemaphoreHandleDesc sem_desc{}; + sem_desc.type = hipExternalSemaphoreHandleTypeD3D12Fence; + sem_desc.handle.win32.handle = fence_handle.handle; + std::printf("STAGE hipImportExternalSemaphore(D3D12Fence)\n"); + hip.check(external.hipImportExternalSemaphore(&imported.semaphore, &sem_desc), "hipImportExternalSemaphore(D3D12Fence)"); + handle_checkpoint("imported_semaphore"); + if (persistent.stream) { imported.stream = persistent.stream; imported.owns_stream = false; } + else hip.check(hip.hipStreamCreateWithFlags(&imported.stream, hipStreamNonBlocking), "hipStreamCreateWithFlags"); + handle_checkpoint("created_stream"); + hip.check(hip.hipMalloc(&imported.errors, sizeof(uint32_t)), "hipMalloc(errors)"); + const uint32_t zero = 0; + hip.check(hip.hipMemcpy(imported.errors, &zero, sizeof(zero), hipMemcpyHostToDevice), "hipMemcpy(error initialization)"); + uint32_t* host_upload = nullptr; + D3D12_RANGE no_read{0, 0}; + dx(upload->Map(0, &no_read, reinterpret_cast(&host_upload)), "Map(upload initialization)"); + for (uint32_t i = 0; i < count; ++i) host_upload[i] = pattern(i, seed); + D3D12_RANGE written{0, static_cast(bytes)}; upload->Unmap(0, &written); + ComPtr allocator; + ComPtr first, second; + dx(device->CreateCommandAllocator(D3D12_COMMAND_LIST_TYPE_DIRECT, IID_PPV_ARGS(allocator.GetAddressOf())), "CreateCommandAllocator"); + dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, + IID_PPV_ARGS(first.GetAddressOf())), "CreateCommandList(input)"); + if (iteration) transition(first.Get(), texture.Get(), D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_COPY_DEST); + linear_copy.pResource = upload.Get(); + first->CopyTextureRegion(&texture_copy, 0, 0, 0, &linear_copy, nullptr); + transition(first.Get(), texture.Get(), D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE); + linear_copy.pResource = shared.Get(); + first->CopyTextureRegion(&linear_copy, 0, 0, 0, &texture_copy, nullptr); + transition(first.Get(), shared.Get(), D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COMMON); + dx(first->Close(), "Close(input)"); + ID3D12CommandList* submit[] = {first.Get()}; queue->ExecuteCommandLists(1, submit); + const UINT64 ready = 2ull * iteration + 1, done = ready + 1; + dx(queue->Signal(shared_fence.Get(), ready), "Signal(D3D12 input ready)"); + hipExternalSemaphoreWaitParams wait{}; wait.params.fence.value = ready; + hip.check(external.hipWaitExternalSemaphoresAsync(&imported.semaphore, &wait, 1, imported.stream), "hipWaitExternalSemaphoresAsync"); + uint32_t n = count, mutable_seed = seed; + void* parameters[] = {&imported.mapped, &n, &mutable_seed, &imported.errors}; + hip.check(hip.hipModuleLaunchKernel(kernel, (n + 127) / 128, 1, 1, 128, 1, 1, 0, + imported.stream, parameters, nullptr), "hipModuleLaunchKernel(shared VRAM)"); + hipExternalSemaphoreSignalParams signal{}; signal.params.fence.value = done; + hip.check(external.hipSignalExternalSemaphoresAsync(&imported.semaphore, &signal, 1, imported.stream), "hipSignalExternalSemaphoresAsync"); + dx(device->CreateCommandList(0, D3D12_COMMAND_LIST_TYPE_DIRECT, allocator.Get(), nullptr, + IID_PPV_ARGS(second.GetAddressOf())), "CreateCommandList(output)"); + transition(second.Get(), shared.Get(), D3D12_RESOURCE_STATE_COMMON, D3D12_RESOURCE_STATE_COPY_SOURCE); + transition(second.Get(), texture.Get(), D3D12_RESOURCE_STATE_COPY_SOURCE, D3D12_RESOURCE_STATE_COPY_DEST); + second->CopyTextureRegion(&texture_copy, 0, 0, 0, &linear_copy, nullptr); + transition(second.Get(), texture.Get(), D3D12_RESOURCE_STATE_COPY_DEST, D3D12_RESOURCE_STATE_COPY_SOURCE); + linear_copy.pResource = readback.Get(); + second->CopyTextureRegion(&linear_copy, 0, 0, 0, &texture_copy, nullptr); + dx(second->Close(), "Close(output)"); + dx(queue->Wait(shared_fence.Get(), done), "Wait(HIP output ready)"); + submit[0] = second.Get(); queue->ExecuteCommandLists(1, submit); + dx(queue->Signal(complete.Get(), iteration + 1), "Signal(completion)"); + dx(complete->SetEventOnCompletion(iteration + 1, event.handle), "SetEventOnCompletion"); + if (WaitForSingleObject(event.handle, 30000) != WAIT_OBJECT_0) throw std::runtime_error("D3D12 fence timeout"); + hip.check(hip.hipStreamSynchronize(imported.stream), "hipStreamSynchronize"); + uint32_t errors = ~0u; + hip.check(hip.hipMemcpy(&errors, imported.errors, sizeof(errors), hipMemcpyDeviceToHost), "hipMemcpy(error verification)"); + if (errors) throw std::runtime_error("HIP input mismatch count=" + std::to_string(errors)); + uint32_t* host_result = nullptr; + D3D12_RANGE read{0, static_cast(bytes)}; + dx(readback->Map(0, &read, reinterpret_cast(&host_result)), "Map(final verification)"); + bool good = true; + for (uint32_t i = 0; i < count; ++i) if (host_result[i] != (pattern(i, seed) ^ 0x3579bdf1u)) { good = false; break; } + readback->Unmap(0, &no_read); + if (!good) throw std::runtime_error("D3D12 output mismatch"); + // Check each release in the successful path; RAII handles all failures. + if (!persistent.stream) hip.check(hip.hipStreamDestroy(imported.stream), "hipStreamDestroy"); + imported.stream = nullptr; + handle_checkpoint("released_stream_reference"); + hip.check(hip.hipFree(imported.errors), "hipFree(errors)"); imported.errors = nullptr; + handle_checkpoint("freed_errors"); + hip.check(hip.hipFree(imported.mapped), "hipFree(external mapping)"); imported.mapped = nullptr; + handle_checkpoint("freed_mapping"); + hip.check(external.hipDestroyExternalMemory(imported.memory), "hipDestroyExternalMemory"); imported.memory = nullptr; + handle_checkpoint("destroyed_memory"); + hip.check(external.hipDestroyExternalSemaphore(imported.semaphore), "hipDestroyExternalSemaphore"); imported.semaphore = nullptr; + handle_checkpoint("destroyed_semaphore"); + shared.Reset(); + if (!CloseHandle(memory_handle.handle)) throw std::runtime_error("CloseHandle(shared memory) failed"); + memory_handle.handle = nullptr; + first.Reset(); second.Reset(); allocator.Reset(); + DWORD cycle_handles = 0; GetProcessHandleCount(GetCurrentProcess(), &cycle_handles); + if (iteration == 0) warmed_handles = cycle_handles; + final_cycle_handles = cycle_handles; + std::printf("CLEANUP iteration=%u handles=%lu warmed_baseline=%lu\n", iteration, cycle_handles, warmed_handles); + std::printf("ROUNDTRIP iteration=%u verified=%u vram_transfer=1 cpu_staging_between_apis=0\n", iteration, count); + } + if (persistent.stream) { + hip.check(hip.hipStreamDestroy(persistent.stream), "hipStreamDestroy(persistent)"); + persistent.stream = nullptr; + } + DWORD handles_after = 0; GetProcessHandleCount(GetCurrentProcess(), &handles_after); + std::printf("HANDLES before=%lu after=%lu\n", handles_before, handles_after); + if (final_cycle_handles > warmed_handles + 4) throw std::runtime_error("Handle count grew after first-cycle warmup"); + std::printf("PASS INTEROP architecture=gfx1201 iterations=%u texture=R32_UINT external_fence=1\n", args.iterations); + return 0; + } catch (const std::exception& e) { + std::fprintf(stderr, "FAIL INTEROP %s\n", e.what()); + loaded_modules(); + return 4; + } +} diff --git a/tools/windows/diagnostic.h b/tools/windows/diagnostic.h index c7b6a1b..78d160c 100644 --- a/tools/windows/diagnostic.h +++ b/tools/windows/diagnostic.h @@ -84,6 +84,8 @@ inline void start() } struct Args { std::string hip_root, cuda_dll, module, context = "primary"; + std::string stream_lifetime = "persistent"; + std::string interop_mode = "roundtrip"; unsigned iterations = 32; int device = -1; Args(int argc, char** argv) { @@ -95,6 +97,8 @@ struct Args { else if (key == "--cuda-dll") cuda_dll = value; else if (key == "--module") module = value; else if (key == "--context") context = value; + else if (key == "--stream-lifetime") stream_lifetime = value; + else if (key == "--interop-mode") interop_mode = value; else if (key == "--iterations") { size_t end = 0; iterations = static_cast(std::stoul(value, &end)); diff --git a/tools/windows/hip_api.h b/tools/windows/hip_api.h index da53175..12ce53c 100644 --- a/tools/windows/hip_api.h +++ b/tools/windows/hip_api.h @@ -27,6 +27,9 @@ struct HipApi { D4R_HIP_FUNCTION(hipModuleGetFunction); D4R_HIP_FUNCTION(hipModuleLaunchKernel); D4R_HIP_FUNCTION(hipDeviceSynchronize); + D4R_HIP_FUNCTION(hipStreamCreateWithFlags); + D4R_HIP_FUNCTION(hipStreamSynchronize); + D4R_HIP_FUNCTION(hipStreamDestroy); #undef D4R_HIP_FUNCTION explicit HipApi(const std::string& root) : search(std::filesystem::path(wide(root)) / L"bin"), diff --git a/tools/windows/hip_stream_lifecycle_probe.cpp b/tools/windows/hip_stream_lifecycle_probe.cpp new file mode 100644 index 0000000..1d14cb6 --- /dev/null +++ b/tools/windows/hip_stream_lifecycle_probe.cpp @@ -0,0 +1,60 @@ +#include "hip_api.h" + +// Isolate Windows HIP stream lifetime from D3D12 and external-resource imports. +int main(int argc, char** argv) +{ + using namespace d4r::diag; + start(); + try { + Args args(argc, argv); + HipApi hip(args.hip_root); + hipDeviceProp_t props{}; + hip.select_gfx1201(args.device, props); + hipModule_t module = nullptr; + hipFunction_t kernel = nullptr; + void* output = nullptr; + struct ResourceCleanup { + HipApi& hip; hipModule_t& module; void*& output; + ~ResourceCleanup() { + if (output) (void)hip.hipFree(output); + if (module) (void)hip.hipModuleUnload(module); + } + } resources{hip, module, output}; + if (!args.module.empty()) { + hip.check(hip.hipModuleLoad(&module, args.module.c_str()), "hipModuleLoad"); + hip.check(hip.hipModuleGetFunction(&kernel, module, "d4r_probe_pattern"), "hipModuleGetFunction"); + hip.check(hip.hipMalloc(&output, 128 * sizeof(uint32_t)), "hipMalloc"); + } + DWORD baseline = 0, count = 0; + for (unsigned i = 0; i < args.iterations; ++i) { + hipStream_t stream = nullptr; + hip.check(hip.hipStreamCreateWithFlags(&stream, hipStreamNonBlocking), "hipStreamCreateWithFlags"); + struct StreamCleanup { + HipApi& hip; hipStream_t& stream; + ~StreamCleanup() { if (stream) { (void)hip.hipStreamSynchronize(stream); (void)hip.hipStreamDestroy(stream); } } + } cleanup{hip, stream}; + if (kernel) { + uint32_t n = 128, seed = i; + void* parameters[] = {&output, &n, &seed}; + hip.check(hip.hipModuleLaunchKernel(kernel, 1, 1, 1, 128, 1, 1, 0, + stream, parameters, nullptr), "hipModuleLaunchKernel"); + } + hip.check(hip.hipStreamSynchronize(stream), "hipStreamSynchronize"); + hip.check(hip.hipStreamDestroy(stream), "hipStreamDestroy"); + stream = nullptr; + if (!GetProcessHandleCount(GetCurrentProcess(), &count)) + throw std::runtime_error("GetProcessHandleCount failed"); + if (i == 1 || (args.iterations == 1 && i == 0)) baseline = count; + std::printf("STREAM_CLEANUP iteration=%u handles=%lu baseline=%lu\n", i, count, baseline); + } + if (output) { hip.check(hip.hipFree(output), "hipFree"); output = nullptr; } + if (module) { hip.check(hip.hipModuleUnload(module), "hipModuleUnload"); module = nullptr; } + if (count > baseline + 4) throw std::runtime_error("HIP stream lifecycle handle growth"); + std::printf("PASS HIP_STREAM iterations=%u handles=%lu\n", args.iterations, count); + return 0; + } catch (const std::exception& e) { + std::fprintf(stderr, "FAIL HIP_STREAM %s\n", e.what()); + loaded_modules(); + return 4; + } +} diff --git a/tools/windows/probe_kernel.hip b/tools/windows/probe_kernel.hip index 9ab0c08..250d25f 100644 --- a/tools/windows/probe_kernel.hip +++ b/tools/windows/probe_kernel.hip @@ -7,3 +7,15 @@ extern "C" __attribute__((global)) void d4r_probe_pattern( if (i < count) output[i] = (i * 1664525u + seed) ^ 0xa5a55a5au; } + +extern "C" __attribute__((global)) void d4r_probe_interop( + unsigned int* shared, unsigned int count, unsigned int seed, unsigned int* errors) +{ + const unsigned int i = __builtin_amdgcn_workgroup_id_x() * 128u + + __builtin_amdgcn_workitem_id_x(); + if (i < count) { + const unsigned int expected = (i * 1664525u + seed) ^ 0xa5a55a5au; + if (shared[i] != expected) __atomic_fetch_add(errors, 1u, __ATOMIC_RELAXED); + shared[i] ^= 0x3579bdf1u; + } +}