ROCm on the Windows lane (-Variant rocm)#
.\windows\Build-Buildkit.ps1 -Variant rocm (amd64 only) builds the AMD image
:winamd64-rocm. Since 2026-09-23 the ROCm layer sits in the sdk slot, like
Dockerfile.nvidia on the GPU lane, so toolchain and media build on top of it and can
turn on ROCm and AMD-GPU features. Every one of them is gated on
(Get-GpuEnvironment).HasRocm: the cpu and nvidia lanes keep their flags and outputs.
Nothing on this page has been seen executing on a GPU. Windows containers get DirectX only, and on this host GPU passthrough is additionally blocked by the 26100/26200 build skew (windows-build-resources.md). Every image check below is GPU-less: presence, linkage, listings, compile-only. GPU proof belongs on the bare host.
What the rocm lane enables#
Component |
Enabled on the rocm lane |
Mechanism |
Needs the ROCm layer |
|---|---|---|---|
GPU API loaders (sdk layer) |
TheRock’s |
OpenCL, Vulkan |
OpenCL: TheRock’s ICD |
GStreamer 1.29.2 |
|
HIP (dlopen), AMF, D3D |
hip: at run time |
FFmpeg n9.0.2 |
AMF encoders, decoders, filters, |
AMF (driver), Vulkan (loader dlopened at run time) |
no |
OpenCV 5.0.0 |
OpenCL T-API (already on every lane); clBLAS/clFFT probes pinned off |
OpenCL |
TheRock’s ICD loader + the registered |
IREE |
HIP HAL driver, |
ROCm/HIP |
at run time |
TVM |
OpenCL runtime; ROCm codegen + runtime as a spike ( |
OpenCL, ROCm/HIP |
spike: yes |
LiteRT-LM |
GPU backend (WebGPU over Dawn on D3D12) |
D3D12 |
no |
ONNX Runtime |
CPU + DirectML, plus the in-tree WebGPU EP as a spike ( |
DirectML, WebGPU (Dawn on D3D12) |
no |
PyTorch |
torch 2.14.1+rocm10.1.0, torchvision 0.29.1+rocm10.1.0, built from source here for gfx1201 and gfx1200 |
ROCm/HIP (AMD’s runtime wheels) |
no |
App venv LiteRT |
|
WebGPU (Dawn on D3D12) |
no |
llama.cpp HIP |
|
ROCm/HIP |
yes: TheRock’s clang, HIP runtime, hipBLAS/rocBLAS |
llama.cpp Vulkan |
the same build’s official Windows Vulkan zip ( |
Vulkan (the sdk layer’s loader) |
no |
MIGraphX + ORT plugin EP |
MIGraphX 2.18.0 from source, |
ROCm/HIP |
yes |
-NoRocmSpikes drops the three spikes: the migraphx stage, TVM_ROCM and ORT_WEBGPU.
Every ONNX Runtime on this lane is the chain’s own build, like on every other lane
(onnxruntime-single-source.md). The WebGPU EP is
built into it, not installed from PyPI.
Not possible on Windows, checked against upstream: the ORT ROCm EP (removed in
1.23), rocDecode / rocJPEG / rocAL (Linux only), any HIP path in FFmpeg, OpenCV or LiteRT,
OpenCV’s MIGraphX dnn backend (NOT UNIX-guarded), IREE’s amdgpu HAL driver (needs
HSA/ROCR), TVM with RCCL.
The ROCm layer (Dockerfile.rocm)#
Built by Build-Buildkit.ps1 -Variant rocm (amd64 only). Since 2026-09-23 the ROCm
layer sits in the sdk slot, where Dockerfile.nvidia sits on the GPU lane (owner
directive 2026-09-23):
base -> sdk = rocm -> toolchain -> media (all branches + merge) -> migraphx -> llama -> torch -> final -> smoke
Toolchain and media now build on top of GPU_TYPE=rocm. That lets media turn on ROCm/AMD-GPU
features behind HasRocm (below). The features stay optional: the cpu and nvidia lanes send the
same configure/build flags as before.
Tags. bk-windows-base is shared. Every tag from sdk on carries a -rocm infix
(Get-BkTag):
bk-windows-sdk-rocmbk-windows-toolchain-rocmbk-windows-media-core-{onnx,ffmpeg,opencv,hailo}-rocm,bk-windows-media-core-rocmbk-windows-media-{litert,tvm}-rocm,bk-windows-media-rocmbk-windows-media-migraphx-rocm,bk-windows-media-llama-rocmbk-windows-torch-rocmbk-winamd64-rocm
A rocm run therefore never overwrites a default image. The cpu and nvidia tags have not moved;
nvidia still shares the default names. A golden table in Driver.Variant.Tests.ps1 locks both
lanes on amd64 and arm64.
The old bk-windows-rocm tag is orphaned: remove it by hand (admin nerdctl --namespace buildkit rmi). Rocm runs from before this change wrote CPU-identical toolchain and media under
the default tags; that is harmless and needs nothing.
Stages and switches.
-Stagestakesbase,sdk,toolchain,media,migraphx,llama,torch,final.migraphxandllamaare rocm-only. Other variants drop them from the default list and refuse them when named.-Stages rocmis gone. It is refused with a pointer to the sdk slot.
migraphxbuildsDockerfile.rocm-migraphx(targetbuilt) FROMbk-windows-media-rocm. Its pins come fromGet-MediaBranchVersionArg -Branch rocm-migraphx.llamabuildsDockerfile.rocm-llama(targetbuilt) FROM the migraphx image.torchbuilds FROM the llama image.-NoRocmSpikesdropsmigraphx, so llama builds FROM the merged media. It also passesTVM_ROCM=0instead of1to media-tvm andORT_WEBGPU=0instead of1to the onnx stage of media-core.These build-args are sent on the rocm lane only (
Get-BkRocmStageArg):TVM_ROCMto media-tvm;ORT_WEBGPUplus the fiveORT_WEBGPU_WINDOWS_*pins tomedia-core-built-onnx(the pins also go under-NoRocmSpikes, where nothing reads them);EXPECT_ROCM_SPIKESto the smoke gate;TORCH_ROCM=1plus everyTORCH_ROCM_WINDOWS_*pin toDockerfile.torch.
The consumers declare them with default
0or valueless, so the cpu and nvidia solves receive nothing new. The rocm sdk stage alone also getsVULKAN_VERSIONandVULKAN_RT_WINDOWS_ZIP_SHA256, and the llama stageLLAMA_CPP_VULKAN_SHA256.-ConcurrentAuxchildren get-Variant rocm(and-NoRocmSpikes). Without it they would build litert/tvm FROM the default toolchain, under default tags.-Variant rocm -Stages torch,finalreusesbk-windows-media-llama-rocm.
Refused at launch:
arm64;
-Gputogether with rocm;a push tag other than
:winamd64-rocm(and a non-rocm run pushing to it);-NoRocmSpikesoff the rocm lane;-Stages migraphxtogether with-NoRocmSpikes;a rocm
-Stageslist that skips a post-media stage between two it builds. Example:media,torchis refused, because torch would build FROM an earlier run’s llama image (backlog #39).
The merge-skip guard also refuses migraphx, llama, torch and final.
Refused in the stage: Install-Rocm.ps1 stops when its base already carries CUDA or ROCm
(GPU_TYPE, CUDA_ROOT, CUDA_PATH). The rocm sdk builds FROM the plain base, never FROM an
nvidia sdk and never on top of a second ROCm layer.
The HasRocm contract. Get-GpuEnvironment returns HasRocm (true only when
GPU_TYPE=rocm) and RocmRoot. It throws when GPU_TYPE=rocm but HIP_PATH/ROCM_PATH has no
lib\cmake\hip, because a mis-plumbed layer would otherwise build CPU flags silently.
Every rocm feature is gated on HasRocm, which the cpu and nvidia lanes can never enter. That
includes features that need the AMD driver but not ROCm (AMF, OpenCL, D3D12, WebGPU), by owner
decision.
ORT builds the cpu lane’s feature set, DirectML included, on this lane and logs ROCm layer present: CPU+DML ORT. The WebGPU spike adds its EP on top (§ ONNX Runtime WebGPU EP).
ORT >= 1.23 has no ROCm EP, and the dead onnxruntime_USE_ROCM branch was deleted.
The app venv’s ORT is checked on every amd64 lane, rocm included. Smoke section 21
requires that C:\opt\OrchestrANT\.venv’s onnxruntime lists DmlExecutionProvider and
that its GenAI reports is_dml_available(). It also requires that the venv’s onnxruntime
is the chain wheel: the only owner of the package, the chain version, and every
.pyd/.dll byte-identical to C:\runtime\wheels\onnxruntime-*.whl. Both are
compiled-in facts, so no device is needed. The bytes matter because PyPI also ships
onnxruntime 1.31.0 for cp314 win_amd64, without DirectML.
CMake isolation. With C:\TheRock\build\bin on PATH, CMake treats C:\TheRock\build as a
package prefix: a PATH entry ending in bin makes its parent a search prefix. TheRock ships
lib/cmake/{flatbuffers,llvm,clang,lld,hip,...} and share/cmake/nlohmann_json there. ORT
v1.31.0 asks for flatbuffers (FIND_PACKAGE_ARGS 23.5.9) and nlohmann_json (3.10) through
find_package first.
Invoke-CmakeConfiguretherefore appends-DCMAKE_IGNORE_PREFIX_PATH=<ROCm tree>on the rocm lane (Get-CMakeRocmIsolationArgs). Elsewhere it appends nothing, so cpu/nvidia command lines are unchanged.A build that needs TheRock’s packages (
find_package(hip), hipblas, …) passes-AllowRocmPrefix.Builds that bypass
Invoke-CmakeConfiguremust isolate themselves or show why they cannot leak. The toolchain’s patched LLVM (raw cmake) cannot leak:its runtimes find LLVM/Clang with
NO_DEFAULT_PATH(llvmorg-23.1.1runtimes/CMakeLists.txt:89-90);TheRock’s zlib/zstd sit only in
lib\rocm_sysdeps;there is no libxml2 and no python.exe in the tree.
A header-only leak (nlohmann_json) leaves no trace in a binary. To check a rocm run, grep its onnx stage log for
TheRock: only the ignore-prefix line should match.
HIP device code compiles only with TheRock’s AMD clang
(C:\TheRock\build\lib\llvm\bin\clang++.exe), passed by absolute path and never through PATH.
The hub’s patched clang-cl is built with LLVM_TARGETS_TO_BUILD=AArch64;X86 and has no AMDGPU
backend, so TVM’s ROCm spike builds an LLVM of its own (§ IREE and TVM).
Host-only HIP consumers stay on clang-cl.
Smoke gate. It runs the CPU suite and floor (170), then Test-RocmImage.ps1 (EXPECT_ROCM=1):
the env contract;
no CUDA;
AMD’s LLVM not shadowing
clang-cl/clang/lld-link;.info\versionmatching the pin;device bitcode present;
a
hipcccompile for gfx1201;every
windows/scripts/build/rocm-checks/*.ps1, sorted by name;then MIGraphX presence and TVM’s
TVM_ROCMmarker matchingEXPECT_ROCM_SPIKES, the mode the driver passed (OrtWebGpu.ps1does the same forORT_WEBGPU).
The rules for the checks:
Each check writes zero or more finding strings (empty = pass) and needs no GPU.
A throw is a finding that names its file.
A missing or empty folder is a finding.
Any finding fails the gate.
Checks run under
Set-StrictMode -Version Latestand$ErrorActionPreference = 'Stop'.Their exit codes are ignored.
The gate bind-mounts
windows/scripts, so a check edit needs no image rebuild.
Source. AMD’s Windows tar install (rocm.docs.amd.com, install → Windows → tar):
https://stable.repo.amd.com/rocm/core/tarball/therock-dist-windows-<family>-<release>.tar.gz.
That is the same TheRock release, from the same host, that
linux/scripts/01-core/setup-rocm-repo.sh installs on Linux as apt packages.
windows/scripts/host/Install-Rocm.ps1 downloads it, verifies it, and extracts it with System32
tar.exe into C:\TheRock\build (AMD’s documented path). It then checks the layout and runs
hipcc --version inside the container. That step now runs on the plain base, with no VS
environment loaded.
No checksum from AMD. Measured 2026-09-22: there is no .sha256, .sig or .asc sidecar,
no SHA256SUMS, and no hash fragment on the pip index. ROCM_WINDOWS_TARBALL_SHA256 is
therefore self-measured and pinned, and the script refuses an empty or malformed pin.
Family |
Covers |
Bytes |
SHA256 |
|---|---|---|---|
|
RDNA4, incl. RX 9070 XT (gfx1201) |
2,244,477,973 |
|
|
every supported GPU |
4,796,456,804 |
|
Disk. The 10.1.0 gfx120X-all tree unpacks to 9.65 GB in 11,575 files (10.0.0: 9.56 GB, 10,544). At its peak the sdk
stage holds the 2.24 GB tarball, the tree and its layer export at once, so Get-StageDiskFloorGb
gives Dockerfile.rocm a 45 GB floor (the default is 40). That figure is calculated, not
measured.
The tree is stored once, as a parent layer, under every later rocm stage, so the toolchain and media floors are unchanged. What grows is the store itself: a second full toolchain + media set beside the default one.
Why the sdk slot, and what it costs (2026-09-23). Until 2026-09-22 the layer forked after
media, because nothing in media consumed ROCm. The owner moved it into the sdk slot, like nvidia,
so media can build ROCm/AMD-GPU features behind HasRocm.
The cost is that the rocm lane shares only base with the default image. Its first run rebuilds
patched LLVM, all of media, the merge and torch. Cold figures from CPU amd64 run 20260919-021220:
patched-llvm 3749 s;
onnx / ffmpeg / opencv / core 1023 / 730 / 859 / 601 s;
litert 3535 s, tvm 2832 s;
merge 4314 s, torch 773 s;
about 5.3 h in total.
sccache hashes no PATH/HIP_PATH/ROCM_PATH for C/C++, so unchanged compile commands may hit
the shared cache; that is unmeasured. A ROCm pin bump, or an edit to Install-Rocm.ps1 or
Dockerfile.rocm, now rebuilds the whole rocm toolchain and media.
The alternative placement, toolchain -> rocm -> media, would keep patched LLVM shared. It was not chosen.
PATH. Dockerfile.rocm sets the variables AMD documents (HIP_PATH, ROCM_PATH,
HIP_PLATFORM, HIP_DEVICE_LIB_PATH, LLVM_PATH) with one deliberate exception: AMD’s page
also adds lib\llvm\bin to PATH. That directory holds AMD’s own clang-cl.exe, clang.exe and
linker, and every compile here must use the patched clang-cl. So only bin\ is added, and it
goes last, because it also carries its own flatc.exe, OpenCL.dll and amdocl64.dll.
Both rules now protect every toolchain and media compile, not just torch.
Rocm.Install.Tests.ps1 fails if either rule, or the Vulkan loader’s placement below, is
broken.
OpenCL ICD and Vulkan loader (sdk layer, since 2026-09-23). Two runtime pieces make the
AMD GPU APIs loadable in the image. Neither needs a GPU to check, and
rocm-checks/GpuLoaders.ps1 checks both.
OpenCL ICD. TheRock’s
bin\OpenCL.dllis the Khronos ICD loader (built from ROCm clr’sopencl/khronos/icd). It loads the vendor ICDs listed underHKLM\SOFTWARE\Khronos\OpenCL\Vendors.Install-Rocm.ps1registersC:\TheRock\build\bin\amdocl64.dllthere, asREG_DWORD 0in the 64-bit view. Until then nothing registered it, so the loader found no platform. The loader rules, read at the rocm-systems commit therock-10.0 pins (6b0e43f):the value name is the DLL path, passed to
LoadLibraryAas is, and only DWORD 0 counts (icd_windows.c:117-151);OCL_ICD_FILENAMESis no substitute: it is read throughsecure_getenv, which returns nothing in an elevated process (icd_windows_envvars.c);ICDs a display adapter registers come first, and the loader dedupes by file name (
icd.c:66-78). With Adrenalin’samdocl64.dllexposed, the driver’s copy wins.
amdocl64.dlllinks PAL statically and imports only system DLLs. AMD’s ICD always reports one platform,AMD Accelerated Parallel Processing, even with no device (clrcl_icd.cpp:119-134). Measured on this host with TheRock’s DLL and no gfx120X GPU enabled: 1 platform, 0 devices, under 1 s. The layout gate now also requiresbin\OpenCL.dllandbin\amdocl64.dll.Vulkan loader. Server Core ships no
vulkan-1.dll, so no Vulkan-linked binary could load in the image. The base image’s Vulkan SDK cannot supply one: scoop installs itcopy_only=1, itsBinhas no loader, andHelpers\VulkanRT.exeis an installer that never runs.windows/scripts/host/Install-VulkanLoader.ps1therefore downloads LunarG’s Runtime Components zip forVULKAN_VERSION:pinned as
VULKAN_RT_WINDOWS_ZIP_SHA256, which equals LunarG’s published digest (https://sdk.lunarg.com/sdk/sha/<v>/windows/<zip>.txt);bump_versions.pyspec_vulkanrefreshes it with the version;the signed x64
vulkan-1.dllgoes intoC:\Windows\System32, whereVulkanRT.exeputs it on a real host. Its version must equalVULKAN_VERSION;the verified copy and
VulkanRT-License.txtgo intoC:\vulkan-loader(VULKAN_LOADER_DIR), which is not on PATH;a System32 copy with other bytes is never overwritten: the build fails instead, so a base image that starts shipping its own loader is noticed.
Why System32 and not a PATH directory: FFmpeg’s Windows
dlopensearches only the application directory and System32 (compat/w32dlfcn.hat n9.0.2), and Python 3.8+ resolves a.pyd’s imports without PATH, which is how TVM’s Vulkan runtime loads. System32 comes before PATH in the standard search, so the copy also wins over any PATH copy; an app-local copy still wins over it. The loader layer comes after the ROCm RUN inDockerfile.rocm, so a Vulkan bump does not re-download the 2.3 GB tarball.The default and nvidia images carry the same loader since 2026-09-26 (BACKLOG CON25), installed by the final stage with
-SystemDir '':C:\vulkan-loaderis appended toPATHand nothing is written to System32. That lets every statically linked consumer load (gstvulkan, IREE’s Vulkan HAL, OpenCV vkcom, ggml-vulkan), not FFmpeg’sdlopenor a.pyd. The rocm variant skips it: itsVULKAN_LOADER_DIRis set.What the loader serves: gstvulkan and amfcodec through it, llama.cpp’s
ggml-vulkan.dll, IREE’s Vulkan HAL, OpenCV’s vkcom backend, FFmpeg’svulkanhwdevice and TVM’s Vulkan runtime. There is no Vulkan ICD in the container, so the loader reports zero devices there. That is the correct answer.Smoke:
rocm-checks/GpuLoaders.ps1. It sorts beforeGStreamer.ps1. It requires thatvulkan-1.dllresolves to System32’s copy and is byte-identical to the pinned one beside its licence, exportsvkGetInstanceProcAddrandvkEnumerateInstanceVersion, and is not older than the SDK’sVK_HEADER_VERSION. It requires the Vendors value (REG_DWORD 0, an existing file),amdocl64.dllloading with its ICD exports, andclGetPlatformIDsreturning 0 or -1001 with an AMD platform listed. Zero platforms is a finding, because AMD’s ICD always reports its platform. Both loaders run in a child pwsh with a 120 s timeout; device counts are only logged.
Proven on 2026-09-25, when the rocm chain’s smoke gate passed (CHANGELOG.md):
GpuLoaders.ps1 found the HKLM value in the final image, and AMD’s ICD loaded there and
listed its platform with no GPU, without a crash or a hang. If PAL ever misbehaves,
dropping the Register-OpenClIcd call is the quick escape; GpuLoaders.ps1 then reports
the missing registration.
Licences. See Redistribution below.
Not proven by anything here: that HIP code runs. Windows containers get DirectX/DirectML only, so a GPU check needs the bare host (Windows 11 25H2, with the RX 9070 XT re-enabled outside any build window):
hipInfo;the family’s
-teststarball;a torch matmul compared against the CPU. This one is required, because upstream TheRock#8379 reports torch+ROCm returning zeros on exactly that card.
HIP compiles against MSVC 14.51#
Under clang, the image’s STL (VS 2026, MSVC 14.51) defines isgreater, isgreaterequal,
isless, islessequal, islessgreater and isunordered as constexpr builtin wrappers. HIP
makes those __host__ __device__, so clang’s __clang_cuda_math_forward_declares.h and
__clang_hip_cmath.h cannot declare their __device__ versions. Every HIP compile in the
image failed: the smoke gate’s one-line hipcc kernel stopped with 20 errors on 2026-09-25, after
215 other assertions had passed.
From then until 2026-10-11 the image shipped a fix inside the toolchain: an overlay of two
headers at C:\runtime\opt\hip-msvc-cmath, each renaming the six names around an #include_next
of the untouched original, loaded by a clang.cfg and clang++.cfg beside TheRock’s clang.exe
(TheRock ships no *.cfg of its own). A probe, Test-HipMsvcCmath.ps1, measured the choice in the
rocm image:
Way to reach TheRock’s clang |
nothing |
|
|
config beside |
|---|---|---|---|---|
|
20 errors |
compiles |
compiles |
compiles |
|
20 errors |
not reached |
compiles |
compiles |
The config files were the only row-complete choice that no user environment could undo.
TheRock 10.1.0 no longer needs it, so the image ships TheRock’s clang as AMD does. With
--no-default-config, which drops the overlay, the probe’s kernel compiles under AMD clang
24.0.0git and MSVC 14.51.36231. This held through hipcc, clang -x hip and amdclang++ -x hip.
First seen on 2026-10-07 with the tarball in a :winamd64 container, then on the published
:winamd64-rocm (32c7cc51…) on 2026-10-11. The overlay, both configs and the probe left the
tree that day (CON78). The first rocm chain without them has to build MIGraphX, ggml-hip and torch.
The smoke gate’s one-line hipcc kernel stays the canary: when a toolset or TheRock bump brings
the clash back, it fails with the same 20 errors, and the overlay comes back from git history.
ONNX Runtime WebGPU EP (rocm lane, spike)#
The rocm lane’s chain ORT (built from source by Build-OnnxFromSource.ps1) is built with
ORT’s in-tree WebGPU execution provider, over Dawn’s D3D12 backend. It replaced the PyPI
onnxruntime-ep-webgpu 0.4.0 plugin the app venv carried for part of 2026-09-23: that
plugin was ORT code built by Microsoft, which the ORT single-source rule forbids
(onnxruntime-single-source.md).
Gating.
It is a spike: on by default with
-Variant rocm, off under-NoRocmSpikes.The driver sends
ORT_WEBGPU=1|0plus the fiveORT_WEBGPU_WINDOWS_*pins tomedia-core-built-onnx(Get-BkRocmStageArg). cpu and nvidia get none of them: the ARGs are valueless, and their configure line is unchanged.ORT_WEBGPU=1anywhere but the native rocm lane is a build error.ORT_ENABLE_WEBGPUis the unrelated Linux knob; the Windows spike does not read it.
Why in-tree (onnxruntime_USE_WEBGPU=ON) and not the plugin build.
The in-tree EP registers inside
onnxruntime.dll, soget_available_providers()listsWebGpuExecutionProviderin every interpreter and GenAI sees it with no extra call.The plugin build (
onnxruntime_USE_EP_API_ADAPTERS) makes a separate provider DLL that each consumer must register withregister_execution_provider_library, and GenAI would not see it without its own registration. The in-tree build needs none of that.ORT’s
onnxruntime_CUSTOM_DAWN_SRC_PATHtakes a pinned, pre-patched Dawn tree and turns Dawn’s own dependency fetch off.
What is fetched, and how each fetch is pinned. Nothing else is fetched. The configure
gate fails if Dawn’s fetch_dawn_dependencies runs.
Input |
Source |
Pin |
|---|---|---|
Dawn at the tag ORT’s |
the GitHub tag archive |
|
ORT’s six Dawn patches |
ORT’s own tree, read from |
ORT’s pin. Applied with GNU |
Dawn DEPS |
as Dawn’s |
the commit read from DEPS; each checkout must equal it |
DXC release ( |
the DirectXShaderCompiler GitHub release |
|
ORT’s other cmake/deps.txt downloads are unchanged: ORT fetches them itself, SHA1-pinned
by ORT. Every rocm onnx build fetches these inputs again; there is no preseed yet.
Why DXC is a pinned release, not built. DXC’s WinIncludes.h includes <atlbase.h>
under _MSC_VER, and the image has no ATL (the same C1083 the LLVM, IREE and TVM builds
record). Its clang-cl branch also adds -Wall, which is -Weverything there, and sets
/Zi in a way sccache cannot cache. So Invoke-DawnPrebuiltDxcPatch replaces Dawn’s
AddSubdirectoryDXC() with an imported dxcompiler target over the pinned release.
DAWN_USE_BUILT_DXC stays ON, so Dawn keeps DXC (ShaderF16, subgroups) instead of falling
back to FXC. The patch refuses to run if that block moves.
Dawn and Tint under clang-cl. Both have upstream clang-cl branches, DAWN_WERROR is off
by default, and no flags or patches are added. That was read from source, and the rocm
chain proved it: the image it built by 2026-09-25 carries the spike and passed
OrtWebGpu.ps1. That was Dawn v20260818.211311; ORT v1.31.0’s v20260916.214021 waits for
the next rocm chain.
What ships.
<ort>\bin\dxcompiler.dll, besideonnxruntime.dll, which loads it at run time. The build refuses anonnxruntime.dllthat imports it ordxil.dll.The same DLL in the chain wheel’s
onnxruntime\capi\, checked byte for byte.No
dxil.dll. ORT v1.31.0 stopped shipping it, because Dawn v20260916.214021 creates the DXIL validator fromdxcompiler.dlland loads nothing else.<ort>\ROCM-FEATURES.txt, the marker:ORT_WEBGPU=0|1and, when on,DAWN_VERSION,DXC_VERSIONandDXCOMPILER_SHA256.
Licences. Dawn and Tint are linked statically into onnxruntime.dll and
onnxruntime_pybind11_state.pyd; their BSD 3-Clause text is already ORT’s own dawn
entry in ThirdPartyNotices.txt. ORT’s notices do not cover DXC, so the build copies
DXC’s three texts to <ort>\licenses\directx-shader-compiler\ and appends them to the
wheel’s onnxruntime\ThirdPartyNotices.txt (Add-OrtWebGpuWheelNotice) before
bdist_wheel. Both are docs/deps/deps.json rows, which Rocm.OrtWebGpu.Tests.ps1
requires.
Build gates.
After configure,
Get-OrtWebGpuConfigureFindingrequiresonnxruntime_USE_WEBGPU=ON,DAWN_FETCH_DEPENDENCIES=OFF,DAWN_USE_BUILT_DXC=ON, D3D12 on, Vulkan off and the prebuilt-DXC log line, and throws otherwise.After the wheel,
Get-OrtWebGpuWheelFindingrequiresWebGpuExecutionProviderin the installed wheel, capi’sdxcompiler.dllequal to the staged hash, nodxil.dllthere, and the DXC notice entry.
Smoke (rocm-checks\OrtWebGpu.ps1, GPU-less). It reads the marker and fails on a mode
that does not match EXPECT_ROCM_SPIKES (a stale parent) or a DLL hash that does not
match. With the spike off it requires no DXC leftovers. With it on, in the base python
and in the app venv, it checks that:
WebGpuExecutionProvideris listed;an identity-model session gets the WebGPU EP, or fails only with
Failed to get a WebGPU adapter;onnxruntime_genailoads that samecapi\onnxruntime.dlland creates a WebGPU model, or fails for want of an adapter;dxcompiler.dllloads from capi and exportsDxcCreateInstance, and nodxil.dllsits in capi or besideonnxruntime.dll.
Any other outcome fails. The whole-script test runs it against a real venv and a shim
python, and each finding carries its [base] or [venv] label.
Proven on 2026-09-25, when the rocm chain’s smoke gate passed: Dawn and Tint compile under the image’s clang-cl, and Server Core without a GPU gets as far as adapter selection. A failure before that has other error text, and the smoke goes red on it by design. Which of the two accepted outcomes it gave is not recorded here.
Not proven yet: Dawn finding the Windows SDK’s d3dcompiler_47.dll in the container,
and the build time the onnx stage gains (probably large).
Bumping. The Dawn pin follows ONNXRUNTIME_VERSION: when ORT moves, take the tag from
its cmake/deps.txt dawn row and re-measure the archive’s SHA256 (held with bump:hold).
DXC is a report row: bump_versions.py (spec_ort_webgpu_dxc) proposes the newest release
and, with --write-all, its single dxc_YYYY_MM_DD.zip name and SHA256.
GStreamer on the rocm lane#
GStreamer 1.29.2 builds its AMD paths on every lane already. Meson auto needs only in-tree stubs, the vendored AMF headers and the Windows SDK. The CPU log out/build-logs/rebuild-amd64-20260919-021219.log installs these with no ROCm input: gsthip.dll + gsthip-0.dll, gstamfcodec.dll, and gstd3d11/gstd3d12 (plugin + -1.0-0 library). So the rocm lane adds pins, isolation and a proof, not new code. All of it is gated on (Get-GpuEnvironment).HasRocm in windows/scripts/build/Build-GstreamerFromSource.ps1. The cpu and nvidia meson lines are byte-identical (SourceBuild.GstreamerRocm.Tests.ps1).
Step |
What |
Why |
|---|---|---|
Pins |
|
|
Isolation |
Entries under the ROCm root are dropped from |
Meson does not configure through |
Proof |
After meson setup, |
A scrub nobody checks is an assumption. |
HIP is loaded at run time, not linked. gst-libs/gst/hip compiles against in-tree stubs. At 1.29.2 the loader opens the first HIP_PATH\bin\amdhip64_*.dll, then amdhip64_7.dll by name. It opens hiprtc as HIP_PATH\bin\hiprtc<MM><mm>.dll, which is hiprtc0716.dll for TheRock 10.1.0 (HIP 7.16; hiprtc0715.dll on 10.0.0). GLib opens both with LoadLibraryW by full path. So amdhip64_7.dll’s own imports, rocm_kpack.dll and amd_comgr.dll, resolve only through PATH or System32, never through HIP_PATH\bin. TheRock’s bin has to stay on the image PATH. Measured against the real 10.0.0 gfx120X-all DLLs, and again against 10.1.0’s on 2026-10-07 (Get-GstHipRuntimeFinding, no finding): amdhip64 exports all 37 names the loader resolves, including hipGLGetDevices/hipGraphicsGLRegisterBuffer, and hiprtc exports all 7.
Not enabled: hip-amd-precompile. At 1.29.2, sys/hip/meson.build looks only for hipcc.bin, and TheRock ships hipcc.exe. Upstream fixed this after the tag (4cdf9d9796, MR 12331). The hiprtc JIT path already covers the converter kernels, so revisit at the next GStreamer bump.
Smoke check: windows/scripts/build/rocm-checks/GStreamer.ps1. It is GPU-less and runs through Test-RocmImage.ps1. It asserts:
the HIP runtime above: exports, import closure, and
hiprtc-builtins<MMmm>.dllonPATH;per plugin (d3d11, d3d12, amfcodec, hip): the DLL exists; its import closure resolves on
PATH/System32; no static link toamdhip64/hiprtc/amfrt; andgst-inspect-1.0 <plugin>loads it, with a private registry and a 300 s timeout.
What it cannot assert:
Elements (
hipupload,hipconvert,amfh264enc,amfh265enc). Without an AMD GPU both plugins register 0 elements and still load: hip finds no device, and amfcodec needsamfrt64.dllfrom the Adrenalin driver plus an AMD adapter.Loading hip in any Windows container, on any lane.
gsthipis built with GL interop, andgstgl-1.0-0.dllimportsOPENGL32.dll, which Server Core does not ship. The load probe is skipped for that reason; everything else is still checked.Loading amfcodec on an image without the Vulkan loader. amfcodec links
gstvulkan, so the same skip applies there. The rocm image shipsvulkan-1.dllin System32 (§ The ROCm layer), so amfcodec IS load-probed on it. At 1.29.2 its Windowsplugin_initis D3D11-only and returns TRUE withoutamfrt64.dll, and gstvulkan registers its elements even whenvkCreateInstancefails for want of an ICD. On the rocm image a missing loader is aGpuLoaders.ps1finding.
Elements have to be proven on the bare host with the RX 9070 XT: gst-inspect-1.0 hipconvert, amfh265enc, d3d12h265dec. Never bake a GPU-less GStreamer registry into an image: these plugins call no gst_plugin_add_dependency.
Not at 1.29.2: onnxinference execution-provider=hip/migraphx/dml. They exist only on GStreamer main (Aug–Sep 2026, untagged).
Evidence: subprojects/gst-plugins-bad/{meson.options,gst-libs/gst/hip/{meson.build,gsthiploader.cpp,gsthiprtc.cpp,gsthip-interop.cpp},sys/hip/{meson.build,plugin.cpp},sys/amfcodec/{meson.build,plugin.cpp,gstamfencoder.cpp},gst-libs/gst/d3d1{1,2}/meson.build,sys/d3d1{1,2}/meson.build} at tag 1.29.2; the TheRock 10.0.0 tarball listing; the CPU build log’s install lines and import-walk summary.
FFmpeg on the rocm lane: AMD AMF and Vulkan#
AMF is below; Vulkan has its own subsection, § Vulkan.
What AMF enables (rocm lane only).
Encoders
h264_amf,hevc_amf,av1_amf.Decoders
h264_amf,hevc_amf,av1_amf,vp9_amf.Filters
vpp_amf(scaling and colour),sr_amf(super resolution),frc_amf(frame-rate conversion), and the screen-capture sourcevsrc_amf(configured asamf_capture).The
amfhwdevice, which can take a D3D11 or D3D12 child device.
The cpu and nvidia lanes are configured exactly as before.
None of it uses ROCm. AMF is AMD’s driver-level media SDK, and FFmpeg n9.0.2 has no HIP code at all (the configure script contains no hip token). It is tied to the rocm lane by owner decision, not by any technical need. The TheRock tarball carries no AMF files, and FFmpeg needs none from it.
How.
Build-FfmpegFromSource.ps1downloads the AMF release’s header-only asset,AMF-headers-<AMF_HEADERS_VERSION>.tar.gz(84 KB, never the 1.2 GB repo). It is verified againstAMF_HEADERS_SHA256inversions.env, which equals the release asset’s GitHub digest.The headers are copied into
compat/amf/AMFin the source tree and configure gets--enable-amf --extra-cflags=-I<src>/compat/amf. This is the same route the ONNX Runtime headers already take on this lane.--enable-amfis explicit, so a missing or too-old header stops configure withamf requested but not found. FFmpeg n9.0.2 requires AMF 1.5.2.0 or newer (configure:7902-7904), and the pin isv1.5.2.After configure, two rocm-only gates fail the stage in minutes rather than at the smoke gate:
ffbuild/config.makmust enable all 12 AMF symbols;ffbuild/config.makmust not name the ROCm tree in any spelling. FFmpeg is not a CMake build, so-DCMAKE_IGNORE_PREFIX_PATHdoes not apply. It is also safe by construction, becauseDockerfile.rocmsets noINCLUDE/LIB/PKG_CONFIG_PATHand no TheRock executable is one of FFmpeg’s build tools.
The headers are also installed to
C:\runtime\ffmpeg\include\AMF. Every lane installslibavutil/hwcontext_amf.h, which includes<AMF/core/Factory.h>; on the rocm lane that header now compiles for consumers.Pins travel as build args:
AMF_HEADERS_VERSIONandAMF_HEADERS_SHA256go through the media-core map inGet-MediaBranchVersionArginto themedia-core-built-ffmpegstage. They are excluded from the merge stage.bump_versions.pyrefreshes the SHA whenever the tag moves.
Smoke check. windows/scripts/build/rocm-checks/FFmpeg.ps1, run by Test-RocmImage.ps1 under EXPECT_ROCM=1. It checks that the image’s ffmpeg lists every AMF encoder, decoder, filter and the amf hwaccel, and that --enable-amf appears in its configuration line. It also checks that the AMF headers are installed and that no amfrt*.dll ships in the image. Since 2026-09-23 it checks the Vulkan components too (§ Vulkan below). Listing reads FFmpeg’s static tables, so no GPU is needed. It does NOT prove that an AMF session opens.
Limits.
Runtime comes from the host driver.
amfrt64.dllbelongs to the AMD Adrenalin driver and is never shipped. FFmpeg’s Windowsdlopensearches only the exe directory and System32 (compat/w32dlfcn.h:134-138), so PATH does not help.Inside a container: UNVERIFIED.
Microsoft supports only DirectX in Windows containers and lists Server Core images as not supported for GPU acceleration.
On this host, container GPU passthrough is blocked by the 26100 image vs 26200 host build skew (docs/windows-build-resources.md).
The only community recipe copies the AMF DLLs from
HostDriverStoreinto System32, and reports no encode test.Verify on the bare host with the exported
C:\runtime\ffmpeginstead:ffmpeg -f lavfi -i testsrc2 -c:v hevc_amf -f null -.
Driver floors: 10-bit encode needs AMF runtime 1.4.32 or newer (driver 23.30); the decoder’s bit-depth detection expects 1.4.36. Per-codec coverage on RX 9070 XT (gfx1201) is UNVERIFIED.
Generic encoder lookup differs per lane. On the rocm lane, looking up an H.264/HEVC encoder by codec id, for example PyAV
add_stream('h264')oravcodec_find_encoder, now returnsh264_amf/hevc_amfinstead of*_d3d12va, because the AMF encoders are registered first (allcodecs.c:864 before 867, 877 before 880). AV1 (av1_d3d12vais registered first) and all decoders are unchanged. The smoke asks formpeg4by name, so it is unaffected.Already on every lane, so not made rocm-only: D3D11VA/D3D12VA/DXVA2 hwaccels, the
*_d3d12vaencoders and the*_d3d11/*_d3d12filters. These come in through autodetect, and making them rocm-only would change the cpu and nvidia flags. The ONNX DNN backend’sdevice=dml(DirectML) is also already built on every lane.Not enabled:
OpenCL: never autodetected. It needs a statically built Khronos ICD loader, because TheRock ships no
OpenCL.liband nocl_d3d11.h.HIP-native FFmpeg: does not exist upstream.
License. The AMF headers are MIT; all 57 files carry the grant, plus AMD’s codec-royalty notice. amf is in no nonfree or GPL list at n9.0.2 (configure:2029-2064, 2192-2195), so the image’s FFmpeg stays --enable-gpl --enable-version3, “GPL version 3 or later”, with no --enable-nonfree.
Evidence.
FFmpeg n9.0.2 at https://raw.githubusercontent.com/FFmpeg/FFmpeg/n9.0.2/ :
configure(350, 2166, 3597, 3609-3730, 4257-4259, 4799-4801, 7902-7904, 8360-8361);libavcodec/allcodecs.c(857-858, 864-865, 877-878, 915);libavfilter/allfilters.c(444-446, 570);libavfilter/vsrc_amf.c(387-388);libavutil/hwcontext.c(91);compat/w32dlfcn.h(134-138).
The AMF v1.5.2 release, https://github.com/GPUOpen-LibrariesAndSDKs/AMF/releases/tag/v1.5.2 : its asset
AMF-headers-v1.5.2.tar.gzhas digest sha256:d3c12eb3…11c99, and the tag peels to commit eadd00804d5f7e5cd8c85d540073198312870776.A real FFmpeg 9.0.2 build with
--enable-amflists exactly these names, and the smoke check reports zero listing findings against it.
Vulkan (rocm lane only, since 2026-09-23)#
What it enables. --enable-vulkan, built against the base image’s Vulkan SDK (VULKAN_VERSION, headers at VK_HEADER_VERSION 363 since 2026-10-07; the list below was measured at 357):
the
vulkanhwdevice, withhwupload/hwmap/hwdownloadto and from it;Vulkan Video decode hwaccels
av1,h264,hevc,vp9, and compute-shader decode hwaccelsapv,dpx,ffv1,prores,prores_raw;encoders
h264_vulkan,hevc_vulkan,av1_vulkan(Vulkan Video) andffv1_vulkan,prores_ks_vulkan(compute);18 filters:
avgblur,blackdetect,blend,bwdif,chromaber,flip,gblur,hflip,interlace,nlmeans,overlay,scale,scdet,transpose,v360,vflip,xfade(each_vulkan), plus the sourcecolor_vulkan;swscale’s Vulkan/SPIR-V backend (
libswscale/vulkan).
vp9_vulkan and av1_vulkan also need vulkan_1_4, that is headers >= 1.4.317 (configure:3547, 3625, 7849). None of it uses ROCm; like AMF it is on this lane by owner decision. The cpu and nvidia configure lines are byte-identical to before, which SourceBuild.FfmpegRocm.Tests.ps1 asserts.
Why it was off. The SDK ships no vulkan.pc, so check_pkg_config_header_only fails. The fallback check_cpp_condition (configure:7806) sees only CFLAGS, and n9.0.2 never assigns $vulkan_incflags. The SDK’s Include missing from INCLUDE was not the cause.
How.
Get-FfmpegVulkanPlanexists only when the AMF plan does, soHasRocmis still read in one place. It fails closed whenVULKAN_SDKis unset or contains whitespace (configure runs the glslc probe unquoted), or whenInclude\vulkan\vulkan.h,Include\spirv-headers\spirv.horBin\glslc.exeis missing.configure gets
--enable-vulkan --extra-cflags=-I<VULKAN_SDK>/Include --glslc=<VULKAN_SDK>/Bin/glslc.exe.--glslcpins the SDK’s compiler; otherwise configure takes the firstglslc,glslangorglslangValidatoron PATH.SPIR-V is compiled at build time; n9.0.2 has no libshaderc or libglslang option. make runs glslc over the 58
*.comp.glslshaders, and FFmpeg’s ownbin2cembeds the result in the libraries (ffbuild/common.mak:113-133). No shader compiler ships.Shader compression follows zlib+gzip autodetection (configure:7275-7285). Which path the image takes is not measured yet; the build logs
config.mak: shader compression ON|OFFright after the Vulkan gate.Build gate.
Get-FfmpegVulkanConfigGapfails the stage unlessffbuild/config.makenables all 35 symbols (CONFIG_VULKAN,CONFIG_VULKAN_1_4,HAVE_SPIRV_HEADERS_SPIRV_H, 9 hwaccels, 5 encoders, 18 filters) andGLSLC=is the plan’s glslc.spirv_compileris in no configure list, so the components glslc builds stand in for it.No new pin. The headers and glslc come from the SDK that
VULKAN_VERSIONalready pins, so aVULKAN_VERSIONbump now also recompiles the rocm FFmpeg.
Run time. FFmpeg never links vulkan-1.dll; it dlopens it (libavutil/hwcontext_vulkan.c load_libvulkan), and its Windows dlopen searches only the application directory and System32. In the rocm image that finds the sdk layer’s System32 copy (§ The ROCm layer); with no Vulkan ICD in a container, device creation then fails. On a bare host it loads the driver’s loader.
Smoke check (rocm-checks/FFmpeg.ps1, no GPU). It requires that -encoders, -filters and -hwaccels list the 5 encoders, the 18 filters and vulkan; that ffmpeg -h decoder=X lists vulkan under “Supported hardware devices” for each of the 9 hwaccel decoders; that the configuration line has --enable-vulkan; and that no DLL or EXE in the FFmpeg bin imports vulkan-1.dll, statically or delay-loaded. An avutil-<major>.dll must be among the files read (hwcontext_vulkan lives there), so the import check cannot pass empty. It does NOT prove that a Vulkan device opens.
Unchanged by it. A generic encoder lookup by codec id still resolves to h264_amf, hevc_amf and av1_d3d12va: the Vulkan encoders register after *_d3d12va and *_amf (allcodecs.c:861, 876, 891). vulkan is in no GPL, nonfree or version3 list (configure:2029-2070, 2192), so the licence stays --enable-gpl --enable-version3. The Vulkan-Headers are Apache-2.0 OR MIT, the SPIRV-Headers MIT, and glslc is a build-time tool only.
Not enabled: libplacebo. FFmpeg’s libplacebo filter needs libplacebo >= 7.351.0 through pkg-config (configure:7425). That would be a new meson source build with its own submodules to pin, plus a shader-compiler library shipped beside FFmpeg. scale_vulkan and the other Vulkan filters already cover GPU scaling.
Limits.
The rocm chain built it with clang-cl on 2026-09-24, once
makedefwas fixed, andFFmpeg.ps1passed on its image on 2026-09-25 (CHANGELOG.md). The earlier host evidence was object files only: a real n9.0.2 configure with Strawberry mingw gcc 13.2 and the same SDK version enabled all 35 symbols, and all 103 Vulkan targets (45 C objects, 58 shaders) compiled. No DLL was linked, and that configure took gcc’s path, not--toolchain=msvc.No GPU-less check can show that a device opens or a shader runs. AMD’s Windows Vulkan Video is reported buggy upstream (AMD-Gfx-Drivers#99; not re-checked). Verify on the bare host with the exported
C:\runtime\ffmpeg, for exampleffmpeg -init_hw_device vulkan=vk:0 -filter_hw_device vk -f lavfi -i testsrc2 -vf format=nv12,hwupload,scale_vulkan=w=1280:h=720,hwdownload,format=nv12 -f null -.
Evidence (n9.0.2): configure 370, 405, 1114-1131, 3192-3677, 4149-4308, 7275-7287, 7800-7861; libavutil/hwcontext_vulkan.c 650-685; compat/w32dlfcn.h; ffbuild/common.mak 113-133; the vulkan/Makefiles under libavcodec, libavfilter and libswscale; fftools/opt_common.c 355-366.
OpenCV on the rocm lane#
What the rocm lane adds: nothing new at build time. The AMD path was already there. OpenCV 5.0.0 has no HIP or ROCm code path. Its GPU paths that work on AMD hardware are all vendor-neutral, and every lane already compiles them:
OpenCL T-API:
WITH_OPENCL=ON,WITH_OPENCL_SVM=ON.D3D11 interop:
WITH_DIRECTX=ON.Vulkan backend of the classic dnn engine:
WITH_VULKAN=ON.
OpenCL builds against OpenCV’s bundled 3rdparty/include/opencl/1.2 headers with no import library. At run time it loads OpenCL.dll by bare name through the standard DLL search and requires clEnqueueReadBufferRect. The rocm layer supplies a Khronos ICD loader, C:\TheRock\build\bin\OpenCL.dll, last on PATH, plus AMD’s PAL-based amdocl64.dll, which Install-Rocm.ps1 registers as an ICD (§ The ROCm layer). That is what gives the T-API a runtime on this lane. Until 2026-09-23 nothing registered the ICD, so the loader found no platform.
Build-OpencvFromSource.ps1 changes only on the rocm lane:
Get-OpencvRocmCmakeArgsadds-DWITH_OPENCLAMDFFT=OFF -DWITH_OPENCLAMDBLAS=OFF. clBLAS and clFFT default ON upstream and are probed through default paths. They are dormant libraries that neither TheRock nor the image ships. On cpu and nvidia the function returns nothing.After configure,
Get-OpencvRocmConfigureFindingfails the stage before the compile in two cases:the summary lacks
OpenCL: YES;any line of
opencv-configure.lognames the ROCm root.
This is the proof that TheRock’s
lib/cmakepackages, zlib or OpenBLAS never reached OpenCV.rocm-checks/OpenCV.ps1, run byTest-RocmImage.ps1, re-checks the shipped bytes and needs no GPU:cv2.getBuildInformation()reportsOpenCL: YESwithLink libraries: Dynamic load;no build-information line names the ROCm tree;
an
OpenCL.dllwith the 1.1 entry point loads through the standard DLL search.
cv2.ocl.haveOpenCL()is reported warn-only, because the smoke container has no GPU. With the ICD registered it should print True, listing AMD’s platform with no device; that is unverified in the image.
Not rocm-specific, but new on 2026-09-23: on every Windows lane OpenCV stopped downloading its own ORT zip, links the chain’s, and ships a working G-API DirectML EP. That change is owned by the ORT page (onnxruntime-single-source.md).
Choosing the device. The host exposes a dGPU and an iGPU, so select one with OPENCV_OPENCL_DEVICE=<platform>:<type>:<name>, for example :GPU:gfx1201. OPENCV_OPENCL_RUNTIME=<path> pins a specific loader.
Not available, and why:
HIP for
cv::cuda/cudevand HIP-in-UMat. Only unmerged upstream PRs exist, validated on Linux only: opencv#29285/#29527 with contrib #4147/#4178, and opencv#29372.MIGraphX dnn backend (opencv#29726). It is
NOT UNIX-guarded, and no Windows MIGraphX exists.RPP HAL (opencv#29505). Closed unmerged.
GPU inside the container. Microsoft supports only DirectX in Windows containers, and on this host the 26100/26200 build skew blocks GPU passthrough. Measure the OpenCL T-API on the bare host or on a 26100 host. See windows-build-resources.md.
Upstream references, all at the 5.0.0 tag: cmake/OpenCVDetectOpenCL.cmake, modules/core/src/opencl/runtime/opencl_core.cpp, CMakeLists.txt:333-344, and TheRock’s win-tarball-list (bin/OpenCL.dll, bin/amdocl64.dll; no clAmd*).
IREE and TVM on the rocm lane#
The ROCm layer (Dockerfile.rocm) sits in the sdk slot, so media-tvm sees GPU_TYPE=rocm and TheRock at C:\TheRock\build. What that turns on:
Component |
rocm lane |
cpu / nvidia |
|---|---|---|
IREE runtime |
HIP HAL driver ( |
unchanged |
IREE compiler |
|
unchanged |
TVM runtime |
|
|
TVM ROCm spike ( |
AMDGPU in a minimal LLVM that TVM builds for itself ( |
off |
IREE needs nothing from TheRock at build time. The HIP headers are vendored (third_party/hip-build-deps, HIP 6.1), so the configure keeps the rocm-lane package-prefix isolation. Two carried details:
The DLL name. IREE’s Windows HIP driver dlopens only
amdhip64.dll, but TheRock 10 and Adrenalin shipamdhip64_7.dll/amdhip64_6.dll. Without a fix the driver reports UNAVAILABLE on every AMD host.Build-IreeFromSource.ps1patches the name list to try_7, then_6, then the legacy name. TheRock’samdhip64_7.dllexports all 72 required symbols (checked 2026-09-23). Upstream candidate: iree-org/ireeruntime/src/iree/hal/drivers/hip/dynamic_symbols.c.The device bitcode. The rocm target downloads ocml/ockl at CONFIGURE time from shark-infra/amdgpu-device-libs v20231101, hash-pinned in IREE’s own CMake.
IREE_ROCM_DEVICE_BC_SHA256in versions.env mirrors that pin (bump:hold, slaved toIREE_VERSION). The build refuses to run when the two differ, or when upstream stops pinning. An IREE bump that moves the archive fails the rocm lane until someone re-derives the pin fromcompiler/plugins/target/ROCM/CMakeLists.txt.
TVM’s find_rocm runs even with USE_ROCM=OFF. It reads $env:ROCM_PATH and puts TheRock’s include\ (flatbuffers, nlohmann, CL, half, thrust, getopt.h) and __HIP_PLATFORM_AMD__ on every translation unit. So without the spike, Build-TvmFromSource.ps1 hides ROCM_PATH for the whole script, configure and scikit-build re-configure alike, then restores it.
The spike carries two patches:
rocm_device_api.cchas no HSA on Windows:kExistaskshipGetDeviceCount.python/tvm/support/rocm.pygets four Windows fixes:ld.lldis found via PATHEXT, thenROCM_PATH\lib\llvm\binby absolute path;the bitcode comes from
HIP_DEVICE_LIB_PATH;the pre-2023
oclc_daz_opt_*andoclc_correctly_rounded_sqrt_*files are optional (TheRock does not ship them, ROCm/TheRock#5448);a missing
rocminfofalls back instead of raising.
Which LLVM TVM links. Every amd64 lane puts the toolchain’s patched LLVM (C:\llvm-patched) first on PATH. It is built with LLVM_TARGETS_TO_BUILD=AArch64;X86 (Build-LlvmFromSource.ps1), so it has no AMDGPU. cpu, nvidia and the rocm lane without the spike link it unchanged, and never query it.
On the spike, Build-TvmFromSource.ps1 (Get-TvmLlvmChoice) asks the PATH llvm-config --targets-built. When AMDGPU is missing, that llvm-config counts as absent, and TVM builds its own minimal LLVM with X86;AArch64;NVPTX;AMDGPU. It uses the same pinned LLVM_WINDOWS_VERSION source tarball, goes through sccache, and installs to C:\temp\llvm-dev\install. That directory is not shipped: the merge copies only C:\runtime\lib\tvm. Adding AMDGPU to the toolchain LLVM instead would re-key every lane.
After the choice, the build reads --targets-built back from the llvm-config TVM will link, and the spike throws when AMDGPU is missing. ROCM-FEATURES.txt records that read-back list as LLVM_TARGETS, not the requested one, so the rocm lane without the spike records AArch64;X86.
Until 2026-09-23 the minimal-LLVM branch ran only when PATH had no llvm-config. On amd64 the spike therefore always found the patched LLVM, threw at the AMDGPU check before configure, and took the whole media-tvm RUN, IREE included, down with it. The rocm chain has since built the minimal LLVM with AMDGPU in the image: its media-tvm stage, IREE and that LLVM included, took 2:55:07, and TVM.ps1 passed (CHANGELOG.md, 2026-09-25). The LLVM’s own share of that time, and its memory use, are unmeasured.
llvm-config must never resolve into the ROCm tree.
Compile only, no GPU. A Windows container gets no HIP compute, so both gates compile and list:
The build gates check that
iree-compiletargets hip gfx1201, that--list_driversshowship, and thatimport tvmloads the OpenCL and ROCm sidecars.The smoke checks
windows/scripts/build/rocm-checks/IREE.ps1andTVM.ps1re-check the image. They assert a linked ELF64 AMDGPU code object withEF_AMDGPU_MACH0x4E (gfx1201) from IREE (CLI and python) and from TVM’s own AMDGPU codegen. They readC:\runtime\lib\tvm\ROCM-FEATURES.txtto know whether the spike ran.TVM.ps1also requires AMDGPU in the marker’sLLVM_TARGETSwhenTVM_ROCM=1. It comparesLLVM_TARGETSboth ways with what the runtime probe reads fromtarget.llvm_get_targets(null when tvm_compiler has no LLVM). TVM names targets byllvm::Triple::getArchTypeName: X86 →x86_64, AArch64 →aarch64, NVPTX →nvptx64, AMDGPU →amdgpuat LLVM 23 (amdgcnbefore the rename, also accepted). A marker that claims a target the compiler lacks, or omits one it links, is a finding.The unit suite runs
TVM.ps1end to end against a fake python and the build’s own marker writer, so dropping either LLVM check, or recording the requested list instead of the read-back one, fails a test.
Running on a host. Kernels execute only on a host with AMD’s driver and an RDNA3/4 GPU; IREE 3.11 rejects gfx1030/gfx1036. Python ≥3.8 does not search PATH for DLLs, so TVM’s ROCm sidecar needs amdhip64_7.dll in System32 (the driver’s copy) or a registered DLL directory.
Not enabled, on purpose:
Feature |
Why |
|---|---|
IREE |
needs HSA/ROCR, which Windows ROCm does not ship |
TVM hipBLAS |
not now |
TVM rocThrust |
forces hipcc/AMD clang as CXX |
TVM RCCL |
not in TheRock 10 for Windows |
Evidence: the research JSON (component “Apache TVM + IREE”). Pinned sources are IREE e4a3b04 (v3.11.0) and TVM 994e0216.
LiteRT-LM GPU backend (rocm lane only)#
What it enables. On -Variant rocm, C:\runtime\lib\litert-lm\bin also carries LiteRT-LM’s Windows GPU runtime, so litert_lm_main.exe --backend=gpu can run:
libLiteRtWebGpuAccelerator.dll(LiteRT’s WebGPU accelerator; a closed-source prebuilt from the Apache-2.0 LiteRT-LM repo);libLiteRtTopKWebGpuSampler.dll;webgpu_dawn.dll(Dawn, which runs on Direct3D 12);DXC’s
dxcompiler.dllanddxil.dll.
DXC’s licence texts go to C:\runtime\lib\litert-lm\licenses\directx-shader-compiler\.
This feature does not use ROCm or HIP. For Windows, Google ships only the WebGPU accelerator: no HIP, Vulkan or OpenCL one (litert/runtime/accelerators/gpu_registry.cc). D3D12 is also the only GPU API that Windows containers accelerate. It is on the rocm lane by owner decision, but it works on any vendor’s GPU.
How. windows/scripts/build/Build-LitertLmBazel.ps1 gates on GPU_TYPE=rocm, the same test as Get-GpuEnvironment’s HasRocm:
[3/6]: LiteRT-LM’s
WORKSPACEmust pin the DXC zip atLITERT_LM_DXC_ZIP_SHA256.[4/6]: the ROCm tree is hidden from Bazel before its server starts. Every env entry under
ROCM_PATH/HIP_PATHis unset or filtered out of list variables (PATH,HIP_PATH,ROCM_PATH,HIP_DEVICE_LIB_PATH,LLVM_PATH, and any list var).HIP_PLATFORMis unset too. The script restores all of it in itsfinallyblock. Bazel does not go throughInvoke-CmakeConfigure, so this is its equivalent ofCMAKE_IGNORE_PREFIX_PATH.[5/6]: one extra target,
@directx_shader_compiler//:dxc_dlls. Bazel fetches it and checks it against LiteRT-LM’s ownWORKSPACEsha256.[6/6]: the script copies these files into
C:\runtime\lib\litert-lm:the three prebuilt DLLs from
C:\llm\prebuilt\windows_x86_64(git-LFS), each checked against its versions.env SHA256 first;the DXC DLLs and licences from
C:\bzl\external\directx_shader_compiler.
On cpu and nvidia, the Bazel command stays build //runtime/engine:litert_lm_main --config=windows --repo_env=ANDROID_NDK_VERSION=, the env is not touched, and no file is added.
Pins. Four keys in linux/scripts/01-core/versions.env, all tied to LITERT_LM_VERSION:
LITERT_LM_WEBGPU_ACCELERATOR_SHA256,LITERT_LM_WEBGPU_SAMPLER_SHA256andLITERT_LM_WEBGPU_DAWN_SHA256: the git-LFS oids ofprebuilt/windows_x86_64/*.dllat the tag;LITERT_LM_DXC_ZIP_SHA256: equal to the GitHub release digest ofdxc_2026_02_20.zipin DXC v1.9.2602.
bump_versions.py (spec_litert_lm, under --write-all) refreshes all four together with the version. They reach the branch through Get-MediaBranchVersionArg media-litert and the ARG+ENV in media-litert-env. They never reach the merge stage.
Smoke. windows/scripts/build/rocm-checks/LiteRtLm.ps1 runs from Test-RocmImage.ps1 with EXPECT_ROCM=1 and needs no GPU. It checks:
the files are present;
each DLL loads in a child pwsh whose PATH is only the bin dir plus System32, and its entry export resolves (
LiteRtAcceleratorImpl,LiteRtTopKWebGpuSampler_Create,wgpuCreateInstance,DxcCreateInstance);--helplists--backend.
Limits.
No GPU run has been proven. GPU in a Windows container needs process isolation plus
--device class/5B45201D-F2F2-4F3B-85BB-30FF1F953599. That passthrough is blocked on this host by the build skew (docs/windows-build-resources.md). Microsoft also lists Server Core base images as unsupported for GPU acceleration.On the bare host, any D3D12 driver (AMD, NVIDIA or Intel) should work; upstream supports that path.
It is unverified whether Dawn accepts the WARP adapter a container can see.
Upstream
--backenddefaults togpu. On the cpu and nvidia images, pass--backend=cpu.Microsoft’s DXC licence terms apply to redistributing
dxil.dllanddxcompiler.dll.
Evidence (all at the pinned versions):
LiteRT-LM v0.17.1:
README.md:141-144; theprebuilt/windows_x86_64tree;WORKSPACE:564-570andBUILD.directx_shader_compiler;runtime/components/sampler_factory.cc:469-482;runtime/engine/BUILD(litert_lm_main exportsLiteRt*viawindows_exported_symbols.def);runtime/engine/litert_lm_main.cc:55(--backenddefaultgpu).LiteRT@9fe5be45 (LiteRT-LM’s pin):
litert/runtime/accelerators/gpu_registry.cc;litert/c/litert_common.h(LITERT_HAS_WEBGPU_SUPPORT_DEFAULT 1).DXC: https://github.com/microsoft/DirectXShaderCompiler/releases/tag/v1.9.2602 (asset digest
sha256:a1e89031…).https://developers.google.com/edge/litert/next/gpu
https://learn.microsoft.com/en-us/virtualization/windowscontainers/deploy-containers/gpu-acceleration
PyTorch on the rocm lane (torch stage)#
The rocm image’s app venv (C:\opt\OrchestrANT\.venv) runs PyTorch built from source in this repository (owner decision 2026-09-29). It has torch 2.14.1+rocm10.1.0 and torchvision 0.29.1+rocm10.1.0 (the image interpreter’s cpXY, win_amd64), built from the upstream v2.14.1 / v0.29.1 tags that PYTORCH_VERSION / TORCHVISION_VERSION name. They are compiled against this image’s ROCm 10.1 SDK (C:\TheRock\build) for both GPUs of the pinned gfx120X-all family: gfx1201 (RX 9070 series) and gfx1200 (RX 9060 series). They load ROCm the way AMD’s wheels do, from AMD’s rocm[libraries] 10.1.0 runtime and a device wheel per GPU. The venv also needs ai-edge-litert (§ The LiteRT extra, below). The cpu and nvidia images are unchanged.
Why here and not AMD’s: AMD’s Windows wheels for ROCm 10.0 stop at torch 2.13.0, and the app locks 2.14.1. torch 2.14 exists for ROCm 7.14 (pytorch.org, Linux only) and for ROCm 10.1 release candidates (rc.repo.amd.com), but not for 10.0.
WebGPU for ONNX Runtime comes from the chain ORT itself (§ ONNX Runtime WebGPU EP). The PyPI plugin onnxruntime-ep-webgpu that this venv carried for part of 2026-09-23 is gone, with its TORCH_ROCM_WINDOWS_ORT_EP_WEBGPU_* pins: it was ORT code not built by this chain.
How it is selected. windows/Dockerfile.torch declares a global ARG TORCH_ROCM=0 and ends with FROM rocm-${TORCH_ROCM}:
appis the old single stage, unchanged:uv syncwithPYTORCH_EXTRA=pytorch-cpu, so the app lock resolves the same way on every lane.rocm-0isappwith no instruction of its own. This is what cpu and nvidia build, with the same cache key and layers as before.torch-rocm-wheelsbuilds the two wheels on the rocmBASE_IMAGE, not onapp, so anAPP_REFmove never re-runs the compile. Onlyrocm-1mounts it, so cpu and nvidia never build it. It has two RUNs, torch and then torchvision. The torchvision ARGs sit below the torch RUN, so a torchvision failure, edit or bump never re-runs the torch compile. One RUN for both cost a finished 68-minute torch build on 2026-09-29, when only torchvision’s import check failed.rocm-1is built only when the driver passesTORCH_ROCM=1. That happens on the rocm lane only (Get-BkRocmStageArg, which also forwards the source build’s versions and the sccache endpoint there).
rocm-1 runs windows/scripts/build/Install-TorchRocm.ps1, then Build-TorchApp.ps1 -Mode verify, so the app’s own smoke suite runs on the ROCm torch. That verify runs the ORT census first: any ONNX Runtime distribution that is not a chain wheel fails the stage (failure-modes.md). The installer itself also requires the venv’s onnxruntime RECORD digest to be the same before and after it runs (Assert-TorchRocmOrtUnchanged).
The source build (windows/scripts/build/Build-TorchRocmFromSource.ps1, then Build-TorchvisionRocmFromSource.ps1, following TheRock’s external-builds/pytorch/build_prod_wheels.py):
Sources.
Save-GitCommitSourcefetchesTORCH_ROCM_WINDOWS_PYTORCH_COMMITand_TORCHVISION_COMMIT, the tags’ commits, and git verifies every object. Submodules come at the superproject’s gitlinks (--depth 1). Each commit’sversion.txt(2.14.0a0upstream) must carry the pinned x.y.z.Build venv on the image’s CPython (uv, no cache): torch’s
requirements.txtandrequirements-build.txt, plusbuild. PyPI’sninjais removed so the image’s 1.13.2 runs; 1.11.1 hangs and 1.13.0 breaks link.exe response files. The pinnedrocmsdist, core and libraries go in too:import torchruns_rocm_init.py, and torchvision’s setup imports torch.The venv gets the base interpreter’s
sitecustomize.pyandpython3.dll, as the app venv does (Build-TorchApp.ps1). The image’s CPython is built with clang, and without that shimsysconfig.get_platform()sayswin32. uv then resolves 32-bit wheels, and the torch wheel would be taggedwin32. The first rocm build of this stage died on exactly that (2026-09-29): a 32-bit Cython failed to load while pyyaml built from its sdist. The builder now stops unless the venv reportswin-amd64.
Before the wheel: HIPIFY (
tools/amd_build/build_amd.py), thentorch/_rocm_init.py(AMD’s content,check_version= the release), force-included into the wheel because scikit-build-core drops files.gitignorenames.OpenMP.
torch_cpu.dllimports MSVC’slibomp140.x86_64.dll(USE_OPENMP, through MSVC’slibomp.lib). TheRock copies it intotorch\lib, but this wheel does not carry it. VS 18 (MSVC 14.51) ships the release DLL only underdebug_nonredist, which Microsoft does not license for redistribution. The image’s VS install puts it inC:\Windows\System32, where the loader finds it, and the builder stops before compiling if it is missing (Assert-TorchRocmSystemLibomp). So the wheel works only in this image, which is the only place it is installed.The environment (
Get-TorchRocmCommonEnv,Get-TorchRocmTorchEnv):ROCM_HOME/ROCM_PATHare the SDK root andCMAKE_PREFIX_PATHis itslib\cmake.TheRock’s
clang-clis CC/CXX, plusHIP_CLANG_PATH.PYTORCH_ROCM_ARCHcomes fromROCM_WINDOWS_GFX_FAMILY, and OpenBLAS from the SDK’slib\host-math.PYTORCH_BUILD_VERSION=2.14.1+rocm<release>(+rocm10.1.0today) andPYTORCH_EXTRA_INSTALL_REQUIREMENTS=rocm[libraries]==<release>.MAX_JOBSallows 5 GB per job, and sccache wraps the host C/C++ when it is configured.
Then
python -m build --wheel --no-isolation, an install into the build venv, and animport torchthat prints the HIP version and the compiled arch list. The torch wheel is staged inC:\torch-rocm-wheels, and the torch tree is removed. The build venv, with torch installed, stays for the next RUN.The torchvision RUN dot-sources the torch builder for its helpers and builds in that venv. It does not use
Start-MigraphxBuildSession, which would reset the work dir. It adds pillow, becauseimport torchvisionimports PIL and torch’s build requirements do not bring it. Thensetup.py bdist_wheel(FORCE_CUDA=1on the GPU-less host) and an import that requires its C++ ops.C:\torch-rocm-wheelsmust then hold exactly the two wheels. Every step logs in full to the persistentC:\sccache-logs(torch-rocm-*.log,torchvision-rocm-*.log).Off in this first build, on purpose:
AOTriton (flash and memory-efficient SDPA,
USE_FLASH_ATTENTION/USE_MEM_EFF_ATTENTION=OFF). Its CMake clones dlfcn-win32, xz and aotriton and pip-installs into a venv of its own at build time, and nothing here pins any of it. SDPA falls back to the math kernel.torch.distributed(gloo/libuv).Image IO in torchvision (PNG, JPEG, WEBP). AMD’s wheels lack it too.
The SDK tree is AMD’s tarball, not the
rocm-sdk-develwheel TheRock’s scripts expand. The tarball holds every devel path at the same place and has no symlinks, so the symlink-privilege expansion bug (TheRock#7807) does not apply. No public precedent for building against the tarball was found. It builds: see the next bullet.Measured 2026-09-29, on TheRock 10.0.0 (
-Variant rocm -Stages torch,finalon the 2026-09-28 rocm parent; the 10.1.0 build is unmeasured until the next rocm image):The torch phase took 4051 s cold and 1719 s with the WebDAV sccache warm from the cold run. 2796 ninja steps,
MAX_JOBS=7.The torchvision RUN took 228 s.
The import printed
2.14.0+rocm10.0.0, HIP7.15.26333, ROCm10.0.0, archsgfx1200 gfx1201.In
rocm-1the app’s wheel smoke passed 13/15 with 0 failures (torch backward, torchvisionops.nms).The smoke gate passed (215 assertions, 1 skipped) and
rocm-checks\Torch.ps1found both GPUs covered.The driver finished in 54 min. The GPU compute limit below still applies.
What the installer does.
Pins. Five files are pinned by URL + SHA256 in
versions.env(TORCH_ROCM_WINDOWS_<NAME>_URL/_SHA256, next toROCM_WINDOWS_*): AMD’srocmsdist,rocm-sdk-core,rocm-sdk-librariesand one device wheel per GPU. They reach the stage as ARG defaults thatsync_versions.pykeeps in step. AMD publishes no hashes and its index carries no#sha256=fragments, so every SHA256 is self-measured (curl -fL <url> | sha256sum).The built wheels (
Get-TorchRocmBuiltWheel): the mountedC:\bkmnt\torch-rocm-wheelsmust hold exactly one torch and one torchvision wheel, each win_amd64 and a+rocm<release>build. Their SHA256 is taken on the spot, so the install stays--require-hashes.Coupling, checked at build time. Every AMD file must be built for
ROCM_WINDOWS_RELEASE. The venv must be the image interpreter’scpXY(PYTHON_VERSION). The app lock’s torch/torchvision must be the release this build made, so a torch move in the app needsPYTORCH_VERSION/TORCHVISION_VERSIONand their commits to follow. Bump the wholeTORCH_ROCM_WINDOWS_*block withROCM_WINDOWS_RELEASE.One device wheel per GPU.
TORCH_ROCM_WINDOWS_SDK_DEVICE[_<GFX>]names the GPU (Get-TorchRocmGpuTarget). The unsuffixed one is gfx1201; gfx1200 uses_GFX1200.rocBLAS coverage. Before any download, every GPU the image’s rocBLAS ships kernels for (
%HIP_PATH%\bin\rocblas\library\TensileLibrary_lazy_<gfx>.dat, the filesrocm-checks\LlamaCpp.ps1reads too) must have a device pin (Assert-TorchRocmGpuCoverage). Pinning more GPUs is allowed. AROCM_WINDOWS_GFX_FAMILYchange therefore fails the torch stage until the device pins follow.Offline install. Pinned files are cached by hash on the torch stage’s uv cache mount (
C:\uvcache\torch-rocm\<sha256>\<file>) and re-hashed before reuse. Each is downloaded to its cache path directly, never as a.partrenamed into place. On the mount that rename failed withCould not find a part of the pathfor the first wheel with a 68-character name, after two shorter ones had moved fine (2026-09-25). That is the create-then-rename class of windows-build-lanes.md § Run-side wcifs symptoms. A partial file is never used: a failed download deletes what it wrote, and every reuse re-hashes. The install isuv pip install --force-reinstall --no-deps --no-index --no-build-isolation --require-hashes. No index is contacted and nothing is resolved. Therocmsdist builds on the venv’s setuptools (83.0.0 from the app lock).The no-GPU trap. The
rocmsdist’ssetup.pycallsoffload-archto choose a GPU family. The build container has no GPU (TheRock’soffload-arch.exesits inlib\llvm\bin, which stays off PATH). The installer setsROCM_SDK_TARGET_FAMILYto the first pinned GPU, gfx1201. That variable takes exactly one GPU, but it only fills the sdist’s genericdeviceextra, and--no-depsnever reads it; the pinned per-GPU wheels stand in for AMD’s per-GPU extras.ROCM_BOOTSTRAP_DISABLE_DETECTION=1is set too. The source build sets the same two for its build venv.Dependency confusion. PyPI has an unrelated
rocm0.1.0. Pinning by URL + hash rules it out.The 10.1.0 index adds
rocm-bootstrapandrocm-sdk-develpackages. The 10.1.0rocmsdist requires neither: against 10.0.0 it only renumbers its pins, drops the Linuxrocm_smi64and adds the Linuxrpplibrary entry (diffed 2026-10-07). TheRock’stherock-10.1Windows preload list equalsGet-TorchRocmInitSource’s, andhiprtc0*.dllmatches the newhiprtc0716.dll.
Smoke. windows/scripts/build/rocm-checks/Torch.ps1 runs under Test-RocmImage.ps1 (EXPECT_ROCM=1) and at the end of the install. It imports torch/torchvision/rocm_sdk without a GPU and asserts:
+rocm<release>builds of torch and torchvisiona non-empty
torch.version.hiptorch.version.rocmandrocm_sdkequal to the releasethe runtime and device packages at the release
torch compiled for every GPU rocBLAS serves (
torch._C._cuda_getArchFlags(), build-time data), and arocm-sdk-device-<gfx>for each (gfx1200 and gfx1201)the chain onnxruntime still imports and lists
DmlExecutionProviderai_edge_litert.interpreterimports, andlibLiteRtWebGpuAccelerator.dllloads withLiteRtAcceleratorImplresolvedevery unconditional requirement of
ai-edge-litertis installed; on the host, the import alone did not catch a missingml_dtypes
It never calls torch.cuda, and torch.cuda.is_available() is False in the container by design.
The LiteRT extra#
ai-edge-litert is LiteRT’s Python package with its prebuilt WebGPU accelerator libLiteRtWebGpuAccelerator.dll (Dawn statically linked), so it works on any vendor’s D3D12 GPU; it is on the rocm lane by owner decision, and rocm-checks\Torch.ps1 requires it.
Until 2026-10-12 it was a PyPI pin, 2.2.0’s cp314 wheel, installed
--no-depsinrocm-1. PyPI has no cp315 wheel (2.3.0, the newest, stops at cp314), so the Python 3.15 pin ended that.It is built from source for the pinned Python (owner decision 2026-10-12, BACKLOG CON91) into the chain’s wheel store, and
Build-TorchApp.ps1installs it on every amd64 variant with the other store wheels. The accelerator stays the 2.3.0 wheel’s closed-source DLL (ML Drift has no public source), pinned asLITERT_WINDOWS_WEBGPU_ACCELERATOR_SHA256and taken fromLITERT_WINDOWS_ACCELERATOR_WHEEL_*.Until that wheel lands, the venv has none:
Build-TorchApp.ps1leaves the lock’s ai-edge-litert out ofuv syncwith a warning when neither the store nor the lock has a wheel for the image’s tag, androcm-1fails onTorch.ps1’s LiteRT check. That red is on purpose; the check is not relaxed for the gap.Its requirements are in the venv: the lock’s own ai-edge-litert dependencies, plus
ml_dtypes, whichBuild-TorchApp.ps1copies in from the base interpreter.Its Dawn falls back to FXC when DXC is absent. Whether Dawn accepts the WARP adapter a container can see is unverified.
Limits.
GPU compute is not proven by the image. The bare-host check (torch matmul on the GPU vs CPU, above) still applies. The upstream TheRock#8379, open, reports zeros from torch on RX 9070 XT with ROCm 7.14.
There is no Windows triton wheel, so GPU
torch.compile/inductor is unavailable.AMD’s runtime wheels carry their own ROCm 10.1.0 (
_rocm_sdk_core/_rocm_sdk_librariesin the venv, about 4 GB installed). It is the same release asC:\TheRock\build, which the source build compiled against, but it is a second copy.No flash or memory-efficient SDPA (AOTriton is off), and no
torch.distributed: § The source build says why.The
torch-rocm-wheelstorch RUN is a one-hour compile (about 30 min with a warm sccache) whenever torch, the SDK or the torch builder moves. A torchvision move re-runs only its own RUN, and an app move neither.The CPU torch installed by
uv syncstays as dead bytes in theapplayer, about 0.5 GB. That is the cost of leaving the cpu/nvidia stage untouched.The rocm lane runs the app verify twice (CPU torch, then ROCm torch), about 1–2 minutes extra.
gfx1200 adds about 430 MB of downloads (the 10.1.0 rocm-sdk device wheel alone is 378,463,821 B) and about 0.7 GB installed.
whl-nextis AMD’s only Windows channel. A withdrawn file fails the download loudly; a cached copy keeps working.
Evidence for AMD’s 2.13 wheels (2026-09-23, Windows 11 host, CPython 3.14.7, uv 0.12.7; the runtime and device wheels are still these, torch and torchvision are now built here):
All ten original files downloaded and hashed; sizes match
Content-Length. The torch wheel is 113,446,960 bytes; rocm-sdk-core is 758,050,981.The gfx1200 wheels were hashed by streaming, sizes equal to
Content-Length(51,106,952, 131,254 and 377,648,520 B). AMD’srocm-10.0.0sdist declares the per-GPU extras (setup.py,src/rocm_sdk/_dist_info.py).The offline install replaced a CPU torch/torchvision in 59 s with the rocm-check clean.
import torchreports2.13.0+rocm10.0.0 hip 7.15.26333 rocm 10.0.0.The app smoke’s CPU ops pass with
HIP_VISIBLE_DEVICES=-1(0 devices).amdhip64_7.dllstatically imports only system DLLs plusamd_comgr/rocm_kpack. The driver is loaded at run time, which is whyimport torchneeds no GPU.
Tests: windows/scripts/tests/Torch.Rocm.Tests.ps1 and, for the source build’s pure parts, Torch.RocmSourceBuild.Tests.ps1.
llama.cpp HIP and Vulkan (Dockerfile.rocm-llama, rocm lane only)#
The stage installs two builds of the same llama.cpp tag, each in its own directory: the HIP build here, built against this image’s ROCm, and upstream’s prebuilt Vulkan build in § The Vulkan build.
What it adds. The rocm image carries llama.cpp b11583 with a HIP backend built here, in C:\runtime\opt\llama.cpp-hip, which LLAMA_CPP_HIP_HOME names. It contains:
ggml-hip.dll, about 129 MB, built from the tag’s source (LLAMA_CPP_HIP_COMMIT, archiveLLAMA_CPP_HIP_SOURCE_SHA256) with TheRock’s AMD clang, for gfx1200 and gfx1201;llama-server.exe,llama-cli.exe,llama.exeand the other tools, and the CPUggml-cpu-*.dllvariants: the tag’s officialllama-b<build>-bin-win-cpu-x64.zip(LLAMA_CPP_CPU_SHA256).
That is how upstream composes its own win-rocm zip: release.yml’s windows-rocm job builds only --target ggml-hip, and the release’s merge step adds the windows-cpu zip’s tools. Two things differ:
No bundled HIP runtime. Upstream ships
amdhip64_7.dll,rocm_kpack.dllandamd_comgr.dllbeside the exes (llama.cpp#26929). Here nothing of ROCm’s sits in the directory:ggml-hip.dllloads the image’s TheRock runtime, hipBLAS and rocBLAS fromC:\TheRock\build\bin, last on PATH.Two GPUs, not twenty.
GPU_TARGETScomes fromROCM_WINDOWS_GFX_FAMILY(gfx120X-all= gfx1200;gfx1201), the GPUs the tarball’s rocBLAS and hipBLASLt have kernels for. Upstream’s DLL carries 20 targets and weighs 973 MB.
Why from source (owner decision 2026-10-07). Upstream still builds against TheRock 10.0.0 (release.yml windows-rocm, ROCM_VERSION: "10.0.0", at b11472 and still at b11583), and every build ships only a win-rocm-10.0 zip. Its ggml-hip.dll and bundled runtime cannot serve a 10.1 image.
Where it sits. base → sdk (rocm) → toolchain → media → [migraphx] → llama → torch → final. The stage is Dockerfile.rocm-llama, target built; Build-Buildkit.ps1 forwards its LLAMA_CPP_* pins, ROCM_WINDOWS_GFX_FAMILY and the sccache endpoint. -NoRocmSpikes builds it directly on the merged media.
How it is built (windows/scripts/build/Build-LlamaCppHipFromSource.ps1):
Rocm lane and amd64 only (
Initialize-MigraphxBuild), and the pins must be well-formed before anything is fetched: a build number and a 40-hex lower-case commit.The archive of
LLAMA_CPP_HIP_COMMITis verified againstLLAMA_CPP_HIP_SOURCE_SHA256(Resolve-PinnedSource/Save-PinnedSource, the MIGraphX helpers).Upstream’s recipe:
-DGGML_HIP=ON -DGGML_BACKEND_DL=ON -DGGML_NATIVE=OFF -DGGML_CPU=OFF, C and C++ by TheRock’sclang.exe/clang++.exepassed by absolute path. On Windows ggml compiles HIP as C++ throughhip::device, so the whole tree uses AMD’s clang; clang-cl has no AMDGPU backend. TheRock 10.1’s HIP headers compile against MSVC 14.51 without the retiredhip-msvc-cmathoverlay (§ HIP compiles against MSVC 14.51).Nothing downloads at build time: the tools, the server’s prebuilt UI (fetched unpinned from a Hugging Face bucket) and BoringSSL are switched off, and
FETCHCONTENT_FULLY_DISCONNECTEDis set.LLAMA_BUILD_NUMBER/_COMMITcome from the pin, because the tarball has no git history.Only the
ggml-hiptarget is built. Its PE machine must be amd64, and the pins it was built from go tollama-cpp-hip-build.jsonnext to it.
How it is installed (windows/scripts/build/Install-LlamaCpp.ps1 -Backend hip -BuiltDir <the build's output>). One installer serves both builds; a backend table (Get-LlamaCppBackendSpec) holds what differs. For HIP it fails closed:
It runs only when
Get-GpuEnvironmentreportsHasRocm, and only when ROCm’s bin hasamdhip64_7.dll,hipblas.dllandrocblas.dll.The build record must name
LLAMA_CPP_HIP_BUILDand the image’sROCM_WINDOWS_RELEASE.The CPU zip is verified against
LLAMA_CPP_CPU_SHA256. It must be flat, hold the load-bearing files (llama-server.exe,llama-cli.exe,ggml-base.dll,ggml.dll,llama.dll), and carry neither aggml-hip.dllnor any DLL name ROCm’s bin has.The tag’s LICENSE is verified against
LLAMA_CPP_HIP_LICENSE_SHA256.
It writes llama-cpp-hip-manifest.json last: size and SHA256 of every file, and the source build’s record. The same RUN then runs rocm-checks\LlamaCpp.ps1 -Backend hip, so a bad pin or a broken build fails this stage.
What the smoke check proves (windows/scripts/build/rocm-checks/LlamaCpp.ps1). Run with no arguments, as the smoke gate runs it, it grades both builds. For HIP it runs with no GPU:
the shipped bytes equal the manifest’s;
the directory is not on PATH;
nothing of ROCm’s sits next to llama-server;
every static import of
ggml-hip.dllresolves, transitively, a ROCm one to ROCm’s bin, and exports what is imported;a child pwsh loads
ggml-hip.dlland resolvesggml_backend_init, withggml-base.dllfrom the llama directory and every ROCm DLL it imports from ROCm’s bin;llama-cli --list-devicesloads ggml-hip and asks the HIP runtime for devices. In a container that answersggml_cuda_init: failed to initialize ROCm: no ROCm-capable device is detectedand exits 0; on a GPU host it lists the devices. Any other outcome, or no such line at all, is a finding;ggml-hip has device code for every GPU that ROCm’s rocBLAS has kernels for;
llama-server --versionreports the pinned build.
Measured 2026-10-07 at b11472 in a :winamd64 container with TheRock 10.1.0 installed by Install-Rocm.ps1, the hub clone mounted:
Build-LlamaCppHipFromSource.ps1builtggml-hip.dll(128,678,400 B) cold in 3 min 27 s, 177 s of it ninja at-j14, beside a MIGraphX build. The log is 19 MB; without-Wno-ignored-attributesit was 160 MB.Install-LlamaCpp.ps1 -Backend hipinstalled it, andLlamaCpp.ps1(each backend, then both as the smoke gate runs it) reported no finding.ggml-hip.dllimportsggml-base.dll,hipblas.dllandamdhip64_7.dll; its device code is gfx1200 and gfx1201.llama-server --version:version: 0.6.0-dev (build 11472, commit d0b490f25),built with Clang 20.1.8, which is the CPU zip’s frontend.llama-cli --list-devicesexits 0 in 0.2 s withggml_cuda_init: failed to initialize ROCm: no ROCm-capable device is detectedandAvailable devices: (none).
Limits.
Nothing proves a kernel runs. Microsoft accelerates only DirectX inside Windows containers, and GPU passthrough on this host is blocked by the 26100/26200 build skew (see windows-build-resources.md).
RDNA4 only.
ggml-hip.dllcarries device code for gfx1200 and gfx1201 alone, the GPUs the gfx120X-all tarball’s rocBLAS/hipBLASLt serve. Any other AMD GPU (for example a gfx1036 iGPU) takes the Vulkan build.On a bare host the driver’s runtime wins. With no bundled copy, the loader takes Adrenalin’s System32
amdhip64_7.dllbefore TheRock’s on PATH (llama.cpp#26929). The image has no System32 copy. Whether a given driver’s runtime serves TheRock 10.1’s hipBLAS is unproven.Driver floor unknown. The minimum Adrenalin driver for TheRock 10.1.0’s HIP 7.16 runtime is not known.
The frontend is upstream’s binary. The tools and CPU backends are the pinned CPU zip, built by upstream with VS’s clang; only
ggml-hip.dllis compiled here. A full source build would lose the server’s web UI and HTTPS, which upstream fetches unpinned at build time.Size. 178 MB in 54 files and +1 layer, down from about 1.24 GB with upstream’s zip.
Pinning. Upstream publishes several builds a day, all flagged prerelease. Renovate reports new builds (approval-gated).
python docs/scripts/bump_versions.py --write-allmoves the build to the newest one that publishes both a win-cpu and a win-vulkan zip, with the tag’s commit, the source archive’s SHA256, both zips’ digests and the LICENSE.
Evidence.
https://github.com/ggml-org/llama.cpp/releases/tag/b11583
https://raw.githubusercontent.com/ggml-org/llama.cpp/b11583/.github/workflows/release.yml (
windows-cpujob, lines 872-939;windows-rocmjob, lines 940-1096)https://raw.githubusercontent.com/ggml-org/llama.cpp/b11583/ggml/src/ggml-hip/CMakeLists.txt (
CXX_IS_HIPCCon WIN32)https://raw.githubusercontent.com/ggml-org/llama.cpp/b11583/tools/ui/CMakeLists.txt (the prebuilt UI download)
https://raw.githubusercontent.com/ggml-org/llama.cpp/b11115/common/arg.cpp (
--versionhandler) and tools/server/server.cpp (llama_serverstart-up order)https://github.com/ggml-org/llama.cpp/issues/26929
The Vulkan build (C:\runtime\opt\llama.cpp-vulkan)#
What it adds. llama.cpp’s official Windows Vulkan zip of the same build, b11583: asset llama-b11583-bin-win-vulkan-x64.zip, 32,112,442 B. It lives in its own directory, which LLAMA_CPP_VULKAN_HOME names, and it is never on PATH. It is the vendor-neutral path to an AMD GPU, and the only one for GPUs the gfx120X-all rocBLAS has no kernels for, such as the gfx1036 iGPU.
Upstream builds only
ggml-vulkan.dll(-DGGML_VULKAN=ON -DGGML_CPU=OFF -DGGML_BACKEND_DL=ON, Vulkan SDK 1.4.357.0) and then adds the windows-cpu zip’s tools.Measured 2026-10-07 at b11472, and again at b11476 and b11583 (2026-10-13): the Vulkan zip is the CPU zip’s 51 files, byte-identical (CRC32 and size), plus
ggml-vulkan.dll. The HIP directory ships the same 51 beside its ownggml-hip.dll.
Pins. Only LLAMA_CPP_VULKAN_SHA256 is Vulkan’s own. The build stays one pin, LLAMA_CPP_HIP_BUILD, the asset name follows from it, and the LICENSE pin is the HIP build’s, because it is the same tag’s LICENSE. The ARG is declared after the HIP RUN, so a Vulkan-only change keeps the HIP layer and its ggml-hip compile cached.
How it is installed: Install-LlamaCpp.ps1 -Backend vulkan, in its own RUN, which then runs rocm-checks\LlamaCpp.ps1 -Backend vulkan. It fails closed: rocm lane only; the zip verified against its SHA256 and flat; ggml-vulkan.dll, ggml-base.dll, ggml.dll, llama.dll and llama-server.exe present; no bundled vulkan-1.dll (the loader comes from the image) and no ROCm DLL; the LICENSE fetched at the tag; llama-cpp-vulkan-manifest.json written last.
The Vulkan loader. ggml-vulkan.dll statically imports vulkan-1.dll (for vkGetInstanceProcAddr, vkGetDeviceProcAddr, vkGetPhysicalDeviceFeatures2 and vkCmdCopyBuffer). In the image that is the sdk layer’s System32 copy (§ The ROCm layer); on a bare host it is Adrenalin’s. The check accepts a loader from System32 or from a PATH entry, and refuses one in the llama directory or the Windows directory.
What the smoke check proves, with no GPU:
the shipped bytes equal the manifest; a file that cannot be read is a finding;
the directory is not on PATH;
a child pwsh loads
ggml-vulkan.dlland resolvesggml_backend_init, withggml-base.dllfrom the llama directory andvulkan-1.dllfrom where the loader order says;vkEnumerateInstanceVersion, which the loader answers without an instance, driver or GPU, reports API 1.2 or later.ggml_vk_instance_initrefuses anything lower (ggml-vulkan.cpp@b11115:4969-4973);llama-server --versionreports the pinned build.
It never calls ggml_backend_init, because that creates a VkInstance. --list-devices is no probe: release builds load backends silently, and ggml_backend_vk_reg returns NULL on any Vulkan error, so a missing loader and a missing GPU both print “(none)”.
Why a separate directory. Each directory is verified file by file against its own manifest. ggml-vulkan.dll beside ggml-hip.dll would load, but ggml_backend_load_all loads HIP first and llama.cpp then skips a second device with the same device_id, so on the RX 9070 XT Vulkan would only add the iGPU. ggml also searches the current directory for backends, so do not run one build with the other’s directory as CWD.
Limits.
No GPU in the container:
ggml-vulkanregisters no device there. On the bare host it uses Adrenalin’s Vulkan ICD. Nothing proves a kernel runs.Size: about +95 MB and +1 layer.
Windows Defender flagged
llama-gguf-split.exe(byte-identical in every zip of a build) asTrojan:Win32/Wacatac.B!ml, an ML heuristic, on this host on 2026-09-23, and quarantined it. The check then reports it missing. The build stores are excluded bySync-DefenderExclusions.ps1; a bare-host export of either llama directory meets the same detection. Whether it is a false positive is unverified.ggml-vulkan compiles in Khronos Vulkan-Headers code, and the zip ships no Khronos licence text. Whether that creates a notice obligation is unverified.
Evidence (at b11115): the release’s asset digest; release.yml 1137-1224 and 1876-1884; ggml/src/ggml-backend-reg.cpp 480-600; ggml/src/ggml-vulkan/ggml-vulkan.cpp 4959-4973 and 15853-15873; src/llama.cpp 222-256.
MIGraphX and the ORT plugin EP (rocm lane, spike)#
-Variant rocm builds windows/Dockerfile.rocm-migraphx (target built) on the merged rocm media, before the llama stage. -NoRocmSpikes skips it. cpu and nvidia never solve it, and no other Dockerfile or media build-arg map carries its pins.
What it produces:
Path |
Contents |
|---|---|
|
AMD MIGraphX 2.18.0 (the head of |
|
AMD’s out-of-tree ONNX Runtime plugin EP |
The MIGraphX stage leaves ORT and GenAI untouched. On this lane they are built with the cpu lane’s feature flags, DirectML included, on the rocm sdk layer and with the rocm-lane CMAKE_IGNORE_PREFIX_PATH. ORT also carries the WebGPU spike, which is on whenever this stage builds (§ ONNX Runtime WebGPU EP). Nobody has compared their bytes with the cpu image’s. The EP is built against that chain ORT and is loaded at run time with onnxruntime.register_execution_provider_library(<name>, r"C:\runtime\lib\onnxruntime-ep-amdgpu\migraphx-ep.dll").
How it is built (Build-MigraphxFromSource.ps1, Build-OrtAmdgpuEpFromSource.ps1, helpers in WindowsMigraphx.Common.psm1):
Two compilers, on purpose. MIGraphX has HIP device code, and the hub’s clang-cl has no AMDGPU backend. MIGraphX therefore compiles with TheRock’s
lib\llvm\bin\clang++.exe, passed by absolute path;lib\llvm\binnever goes on PATH. Its host-only deps (abseil 20250512.0, protobuf v30.0, msgpack-c cpp-3.3.0, SQLite 3.50.4) build with clang-cl, as upstream’s Windows CI does. The EP is host-only C++ and builds with clang-cl.Targets are explicit.
GPU_TARGETScomes fromROCM_WINDOWS_GFX_FAMILY(gfx120X-all=gfx1200;gfx1201), never from a host probe.Configuration no upstream CI builds. rocMLIR is not in the TheRock tarball, so
MIGRAPHX_ENABLE_MLIR=OFF. That costs the rocMLIR fusions (performance, not correctness). Also used:BUILD_DEV=OFF(hiprtc JIT), composable_kernel off, CMake 4.4.3, TheRock 10.1. Upstream’s only Windows CI is build-only: gfx942,BUILD_DEV=On, ROCm 7.13.CRT. MIGraphX and its deps are
/MD. The EP is/MT, as upstream forces it; it talks to ORT and MIGraphX only through C APIs.Supply chain. Every archive is SHA256-pinned in
versions.env(MIGRAPHX_WINDOWS_*,ORT_AMDGPU_EP_*). The EP declares 8 FetchContent URLs (11 with DirectML) with noURL_HASH. Each one is pre-seeded throughFETCHCONTENT_SOURCE_DIR_*from a verified archive, and so is the abseil that protobuf v34.1 fetches. The build refuses when the pinned commit declares a URL no pin rebuilds byte-for-byte, andFETCHCONTENT_FULLY_DISCONNECTEDwith CMP0170 turns any other fetch into a configure error. rocm-cmake is the one dependency with no archive pin. MIGraphXrocm-10.0callsrocm_add_version_resource(rocm-cmake 33541cd51f), and TheRock’s rocm-cmake predates it by four commits:Unknown CMake command, 2026-09-25. So phase 2 reads the commit that MIGraphX’s ownrequirements.txtpins, from the SHA256-pinned MIGraphX tree (Get-MigraphxRocmCmakeCommit, 40-hex only). It fetches exactly that commit with git, which verifies every object against the id (Save-GitCommitSource), and installs it into the deps prefix, which precedes TheRock onCMAKE_PREFIX_PATH. A MIGraphX bump carries its own rocm-cmake, with nothing to re-derive inversions.env. nlohmann_json stays TheRock’s, through a two-file package shim in the deps prefix (Write-NlohmannJsonConfigShim). TheRock’s copy was installed by an MSVC-style build, so its exported target lists<prefix>/nlohmann_json.natvisas an interface source, and TheRock’s dist does not ship that file (Cannot find source file, 2026-09-25). The shim includes TheRock’s config and clears that property. Until 2026-10-11 HIP sources compiled with a two-header overlay (§ HIP compiles against MSVC 14.51), at first passed by MIGraphX’s own-isystem(Write-HipMsvcCmathOverlay). Under clang, MSVC 14.51’s<cmath>definesisgreater,isgreaterequal,isless,islessequal,islessgreaterandisunorderedasconstexprbuiltin wrappers, which HIP makes__host__ __device__. clang’s__clang_cuda_math_forward_declares.hand__clang_hip_cmath.hthen cannot declare their__device__versions (114 errors, 2026-09-25). Each overlay header renamed those six names around an#include_nextof the untouched original, so device code used MSVC’s versions; TheRock 10.1 needs none of it.-Xclang -fno-cuda-host-device-constexprfixes these errors but breaks device code, becausestd::integral_constant’s conversion becomes host-only. A VS 2022 toolset was declined in favour of the image’s own toolchain. MIGraphX builds unpatched since 2.18.0: the pin carries upstream 5a80dc91ba. Atrocm-10.0an MLIR-off build leftis_module_fusible,adjust_param_shapes,dump_mlir_to_fileanddump_mlir_to_mxrundefined, andmigraphx_gpu.dllfailed to link (2026-09-25);patches\migraphx\001-mlir-off-stubs.patchbackported the fix until 2026-10-07. Every pinned archive is extracted byExpand-PinnedArchive, which skips the symlinks that 7-Zip 25+ refuses as dangerous: any link whose target climbs with.., even one that stays inside the tree. The flatbuffers seed carries nine (Java test dirs,ts/package.json), and no CMake file reads them. Any other 7-Zip error still fails the stage (2026-09-25). It also deletes the intermediate.tarof a.tar.gz. The EP’s own source leavesonnxruntime-ep-amdgpu.tar, which is an ONNX Runtime archive by G2’s name rule, so the EP’s verify phase refused it.Isolation. The deps configure with the rocm-lane
CMAKE_IGNORE_PREFIX_PATH. MIGraphX and the EP pass-AllowRocmPrefix, because they need hip, MIOpen, rocBLAS, hipBLASLt and hiprtc from TheRock.Chain ORT only. The EP build ends in the ORT build gate (
Assert-ChainOrtOnly -Consumer amdgpu-ep, phase 5), which reads its configure log,CMakeCache.txtandbuild.ninjaand writes the stamp the image census requires (onnxruntime-single-source.md).Measured 2026-10-07, 2.18.0 on TheRock 10.1.0, in a
:winamd64container with the tarball installed byInstall-Rocm.ps1: the build ran in 58 min (host deps 128.5 s, configure 76 s, build and install 3260 s), andmigraphx_gpu.dlllinked unpatched with MLIR off.src/targets/gpu/prefuse_ops.cpp, a host-only file unchanged sincerocm-10.0, took about 50 min of that on one core at-O3with AMD clang 24.0.0git, so it is the stage’s critical path. At-O0it still ran past 11 min, so the cost is in the front end, not the optimizer. The EP then built against it in 7 min 19 s, androcm-checks\MIGraphX.ps1found nothing:migraphx-ep.dllloads GPU-less and exportsCreateEpFactories.
Smoke (windows/scripts/build/rocm-checks/MIGraphX.ps1, GPU-less) checks:
both trees are complete;
every static import of a MIGraphX-built PE resolves beside it, in
HIP_PATH\binor in System32;the TheRock HIP runtime set sits beside the EP, byte-identical;
amdhip64is only delay-imported;migraphx-ep.dllloads (LoadLibraryEx) and exportsCreateEpFactorieswithout starting HIP.
It does not call RegisterExecutionProviderLibrary. The EP’s GetSupportedDevices calls hipGetDeviceCount() first, so registration needs a HIP device (mgx_factory.cc), and Windows containers get DirectX only.
Limits:
Run it on the bare host. Put
%HIP_PATH%\binon PATH, or copy the TheRock tree. The EP directory holds the load closure; MIOpen, rocBLAS, hipBLASLt androcm_kpack.dllwith their kernel payloads stay in TheRock’sbin.Prove it against CPU output. Compare against ORT’s CPU EP (TheRock#8379: zeros on gfx1201).
Redistribution is open. The bundled
amdhip64_7.dllcarries the same open redistribution question as the rest of the TheRock tree (see Redistribution).GenAI is not wired. onnxruntime-genai’s AMDGPU EP (v0.16.0+) needs the
amdgpu-ep.dllumbrella (-DUSE_AMDGPU=ON, which renames the MIGraphX backend tomigraphx-backend.dll). It is not built here.
Evidence:
AMDMIGraphX CMakeLists/requirements/windows.yaml at the pinned 2.18.0 commit; its
requirements.txtdiffers fromrocm-10.0’s only in the rocMLIR line, which an MLIR-off build never readssrc/migraphx/CMakeLists.txt (/DELAYLOAD, POST_BUILD closure)
src/shared/CMakeLists.txt (HIP DLL list without amd_comgr.dll)
Redistribution#
The TheRock tree ships licence notices for its libraries under share\doc: MIT (rocBLAS,
hipBLASLt, MIOpen, rocFFT, rocRAND, rocSPARSE, hipcc, …), BSD (rocSOLVER, hipCUB) and
Apache-2.0 (rocThrust; LLVM with the LLVM exception). The HIP runtime (amdhip64_7.dll,
hiprtc*) and the OpenCL runtime (OpenCL.dll, amdocl64.dll) ship without a
licence file, and the Windows HIP runtime links AMD’s prebuilt PAL. The PyTorch wheels
carry copies of the same runtime; neither llama.cpp directory carries one (since b11472 the
HIP build loads the image’s).
The owner cleared that question: :winamd64-rocm was published on 2026-09-28 (image
built at ad08bc30), with the runtime shipped as TheRock delivers it. docs/deps/deps.json
keeps LicenseRef-Proprietary-EULA for the runtime: the id describes the licence, not the
permission.
The Vulkan loader is LunarG’s signed build of the Khronos Vulkan-Loader, Apache-2.0 with
MIT parts; its VulkanRT-License.txt ships in C:\vulkan-loader. The DXC dxcompiler.dll beside the
chain ORT carries Microsoft’s licence terms (LICENSE-MS.txt) next to NCSA, the same
review question as the LiteRT-LM copy. ai-edge-litert’s wheel ships no licence or NOTICE
text of its own.
Open points#
HIP inside a Windows container is unproven and cannot be measured on this host (build skew). GStreamer’s
gsthipcannot load in any Server Core container on any lane at all: it linksgstgl, which importsOPENGL32.dll.GPU correctness: upstream TheRock#8379 reports torch+ROCm returning zeros on the RX 9070 XT. Compare every GPU path against its CPU result on the bare host.
The spikes build, but none has run. By 2026-09-25 all three built in the image and passed their GPU-less smoke checks: MIGraphX with
MLIR=OFF,BUILD_DEV=OFF, gfx120X and TheRock 10 (a configuration no upstream CI builds; one backported patch until 2.18.0), TVM’s ROCm runtime on its minimal LLVM with AMDGPU, and Dawn and Tint under the image’s clang-cl. No GPU has executed any of them.The 2026-09-23 additions have run in a container. By 2026-09-25 the OpenCL ICD registration, the Vulkan loader, FFmpeg’s Vulkan build, the llama.cpp Vulkan zip, gfx1200 and
ai-edge-litertwere in the rocm image, and their GPU-less checks passed (GpuLoaders.ps1,FFmpeg.ps1,LlamaCpp.ps1,Torch.ps1). Each section says what its check cannot see.Found on every lane, not fixed here (owner decision): LiteRT-LM is cloned by tag, not by commit. OpenCV’s configure-time ORT download, listed here until 2026-09-23, was a Windows-lane defect (Linux pre-set
HAVE_ONNXRUNTIMEall along) and is fixed on every Windows lane (onnxruntime-single-source.md).