feat(windows): enable native HIP with deterministic runtime binding - #788
feat(windows): enable native HIP with deterministic runtime binding#788Kenneth-Javier wants to merge 15 commits into
Conversation
|
CI approved and running — it was stuck in Read the description properly and there is a lot here I like, so let me say what stands out before the review itself. You reused the existing seam instead of inventing one. Fail-closed runtime binding. Refusing to start the GPU tier when a different 78 loader-contract tests with synthetic stub DLLs, running on CI without a GPU. That is the part I would not have thought to ask for. Most GPU PRs here can only be verified by whoever owns the silicon; you made most of yours checkable by anyone. What I will be reading forNot a rubber stamp — 4,273 lines touching the build and the loader:
One questionYou validated on physical If it is token-exact, say so in the PR body — it is the strongest claim you can make and it is currently not stated. |
|
CI ran (it was held in The loader joins with Worth noting the two tests that failed are exactly the ones checking that a missing runtime is reported and that identity resolution is sound — the security-relevant pair. Good place for the harness to be strict; it just needs to be strict about the right thing. Also now CONFLICTING — five PRs merged to Everything I said in the earlier comment stands, and the failure does not touch any of it: the fail-closed binding, the reuse of the existing The question from before is still the one I care most about: was the teacher-forced |
|
Yes — it was token-exact against the CPU path at the token-ID level. In the retained authoritative sweep, CPU teacher forcing ran twice and hybrid This was not a close or tolerance-based comparison. I am not claiming Separately, the free autoregressive decode also produced the same exact I updated the PR body to make the teacher-forcing token-exactness explicit. |
57145e5 to
e5ff643
Compare
Post-rebase validationAfter the physical validation above, upstream The only manual conflict resolution was in Post-rebase validation on head
The physical |
Add a portable native Windows build path for the HIP GPU backend while keeping the default CPU build and the host-only HIP_DLL mode independent of the HIP SDK. - build backend_cuda.cu as coli_hip.dll for explicit gfx targets - select the SDK through HIP_SDK_ROOT with component-level overrides - preserve the existing coli_cuda_* backend ABI - keep HIP_DLL host builds free of amdhip64 and HIP SDK requirements - adapt TEMP validation to the MSVC CRT - clean and ignore generated HIP DLL artifacts - document the experimental build workflow and runtime limitations - repair the unmatched Windows documentation fence so GPU sections render Runtime loading of coli_hip.dll is intentionally deferred. The current loader still selects coli_cuda.dll, and deterministic amdhip64_7.dll resolution is part of the subsequent runtime-loader work. Validated on Windows 11 with Radeon 8060S gfx1151, a Windows HIP SDK and MSVC 14.44: - PE32+ x64 DLL with embedded gfx1151 device code - exact 48-to-48 export and loader-contract equality - only amdhip64_7.dll and KERNEL32.dll imported - no embedded absolute SDK path - 36 C test binaries passed - 273 Python tests passed, with 40 skipped Windows HIP does not currently have hosted CI coverage.
Select the Windows backend DLL name and loader diagnostic label from COLI_HIP_DLL so HIP hosts seek coli_hip.dll while CUDA hosts retain their existing behaviour. - derive bounds, copying and both load attempts from one filename constant - emit HIP-specific loader diagnostics in HIP mode - preserve byte-identical CUDA loader diagnostics - retain the shared coli_cuda_* backend ABI - strengthen the existing CUDA environment test owner with CPU-only subprocess coverage for both DLL modes Deterministic amdhip64_7.dll selection and runtime-path verification remain deferred to the next loader slice.
The engine enters its CUDA block only for a truthy COLI_CUDA, so an unset variable means the GPU was never requested: no backend load, no diagnostic, CPU path. Two tests claimed otherwise and both were unreachable rather than merely stale. One asserted "auto-detect: backend unavailable", a string git history shows was introduced by the test itself and has never existed in production. The other asserted "[CUDA] mode:" with COLI_CUDA unset, although that line is emitted only when the backend was explicitly enabled. The Windows auto-enable they described belongs to the separate coli launcher, which sets COLI_CUDA=1 before starting the engine. - state the opt-in contract instead of auto-detect throughout the module - scrub COLI_CUDA, COLI_GPU and COLI_GPUS from child environments - assert no CUDA diagnostic, loader miss or mode line on the unset path - prove execution reaches model validation instead of failing in GPU setup - correct the GPU-present scenario so hardware alone does not opt in - add one non-skipping isolated CUDA_DLL-host contract test Both historical tests still skip without a CUDA host, so the isolated host test carries the contract on machines that have none. No production behaviour changes.
test_cuda_env.py had grown to cover two unrelated contracts: the COLI_CUDA environment semantics it was written for and the backend-DLL selection added for Windows HIP. Split the second into its own module so each file states one thing and the loader tests no longer pay for the other module's import-time probe of the resident colibri binary. - move LoaderBackendSelectionTest to test_backend_loader.py - preserve every test name, executable code line and assertion - retain isolated CUDA_DLL and HIP_DLL host builds and sandboxes - keep test_cuda_env.py focused on COLI_CUDA startup semantics - use private fixture helpers instead of importing a module with import-time binary and GPU probes - remove the now-unused shutil import from test_cuda_env.py No production behaviour or test semantics change. The suite remains at 277 Python tests with 40 skipped, split as seven environment tests and four loader tests for the extracted owner.
Add CPU-only Windows DLL fixtures that model two distinguishable HIP runtimes and one complete fake HIP backend without requiring ROCm, a HIP SDK or GPU hardware. - derive the current 48-symbol backend ABI directly from backend_loader.c - build runtime A and runtime B with the production amdhip64_7.dll basename - link a fake coli_hip.dll against the runtime marker import - export the complete coli_cuda_* ABI plus one test-only bound-runtime accessor - inspect PE architecture, imports and exports with the existing MSYS2 objdump - build everything under unique temporary paths containing spaces - verify explicit cleanup leaves no generated source or binary in the repo The fixtures are structural in this slice: no DLL is loaded and no behavioural runtime-marker claim is made. A native same-process harness will exercise the binding contract separately. No production behaviour changes and no new toolchain installation is required.
Add a native Windows test harness that compiles the real backend_loader.c unchanged and exercises it in a controlled subprocess with distinguishable amdhip64_7.dll stubs. - preload runtime A or runtime B from an absolute Unicode path - inventory loaded same-basename runtime modules at each loader stage - verify the preload marker, loaded path and backend-bound marker agree - prove an already loaded runtime B satisfies a backend linked against A - classify the ambient no-preload result without assuming System32 behaviour - keep all DLL loading outside the Python unittest process - extend explicit fixture cleanup to include the native harness artifacts The tests demonstrate that same-process preload silently controls backend dependency binding. This provides the behavioural contract required for a future fail-closed production preload and path-verification implementation. No production behaviour changes and no GPU operation is performed.
Require COLI_HIP_RUNTIME_DIR when a Windows HIP host actually enters the backend loader. Read and validate the setting with Unicode Windows APIs before the existing backend-DLL load is attempted. - distinguish a missing variable from an explicitly empty value - reject relative and drive-relative paths - require the configured path to exist as a directory - require amdhip64_7.dll to exist beneath it as a file - support spaces, Unicode values and trailing separators - leave CUDA-mode loading completely unaffected - cover the contract with CPU-only native-harness tests This commit validates configuration only. It does not preload the runtime, inspect loaded modules or enforce that the configured runtime supplies the backend dependency. The existing configured-A and preloaded-B witness records that unsafe behaviour for the subsequent binding-enforcement change.
Preserve and report the Windows errors from both existing backend-DLL load attempts instead of describing every failure as a missing backend file. - distinguish an absent candidate from an existing file that failed to load - retain separate primary and fallback Windows error codes - classify missing modules or dependencies, missing procedures and invalid PEs - always include the numeric Windows loader error - avoid unsupported claims about specific dependencies or ambient modules - cover each diagnostic class with deterministic CPU-only DLL fixtures The backend names, load order, search flags and fallback behaviour are unchanged. This commit improves diagnostics only; it does not preload a HIP runtime, inspect loaded modules or enforce runtime binding.
Add the Windows loader primitives needed to reason about a configured HIP runtime before the production loader begins enforcing that decision. - grow Unicode module and final-path buffers without a MAX_PATH bound - compare physical files through symmetric modern or legacy identities - enumerate every loaded amdhip64_7.dll with Toolhelp32 - distinguish empty, incomplete, single and multiple runtime inventories - reduce verified inventory state to an explicit five-way binding decision - cover hardlink aliases, repeated references and simultaneous runtimes - expose value-only helper wrappers only in dedicated native test builds The helpers are not called from coli_cuda_load and do not preload, accept or reject a runtime. Runtime enforcement, Unicode HIP backend loading, fallback removal and post-load verification remain for W1-B2d2.
Require a Windows HIP host to prove that the configured amdhip64_7.dll is the only matching runtime before the backend may become available. - inventory every loaded same-name runtime and compare physical file identity - fail closed on wrong, multiple or unverifiable runtime state - acquire one exact-path runtime reference and verify it after mapping - load coli_hip.dll only from the absolute Unicode executable directory - verify the runtime again after the backend maps and before symbol resolution - release owned runtime and backend references on every failed initialization - preserve the existing CUDA loader and document the experimental HIP setup The post-load identity check detects path replacement after mapping but does not make path-based loading atomic. Real gfx1151 runtime and model validation remain for W1-B3.
Record how to build the synthetic glm_tiny oracle model without touching the system Python or the repository. - mark the tools requirements as development-only, not a Colibri runtime need - name the minimal direct set for make_glm_oracle.py: torch, transformers, safetensors - explain why transformers>=5.11 is a hard floor rather than a preference - add an isolated CPU-only environment recipe and a verification command - document that the generator has no --outdir and writes relative to its working directory, so it must run outside the repository - show the teacher-forcing self-test against externally generated files The model itself is not generated here, and no token has been verified.
The Windows HIP DLL path now builds AND loads: the host binds the runtime named by COLI_HIP_RUNTIME_DIR and fails closed on an identity or duplicate-runtime mismatch, so "build support only" and the deferred-runtime-selection note are both stale. Replace the Limitations list with what was actually tested and what it does not claim: one configuration (Radeon 8060S reporting gfx1151, TheRock HIP 7.14.60850, MSVC 14.44.35207, Windows SDK 10.0.26100.0); dense tensors on HIP with routed experts still on CPU, which is not full-GPU MoE; the device line proves nothing without "resident set: N tensors"; per-tensor CPU fallback keeps the exit code 0; and normal model exits still rely on process teardown, as they already do on the existing host path. Still no hosted CI for the Windows HIP build or runtime - no runner provides the toolchain or a GPU. Detailed setup stays in docs/windows.md.
…residency check The AMD section still said no end-to-end GPU inference had been validated. It has been, on one configuration, so state that precisely and fence off everything it does not cover. Add the three things that actually block a first successful build: - pin VCToolsVersion to 14.44.35207. hipcc's clang picks the newest installed MSVC toolset, and 14.51 (Visual Studio 2026) has a <cmath> that collides with HIP __device__ overloads. - override a machine-wide HIP_PATH inside the build shell only, and clear the ROCm toolchain variables, so a second install cannot redirect headers or device bitcode. - System32 on PATH and a writable TMP/TEMP/TMPDIR, both of which MSYS2 shells commonly leave unset. Add how to put tensors on the GPU and how to check that it happened: COLI_CUDA=1 alone places nothing, CUDA_DENSE=1 makes the dense tensors eligible, routed experts stay on CPU, and "resident set: N tensors" - not the device line - is the evidence. Document that per-tensor upload failure falls back to CPU with exit 0, and that normal model exits rely on process teardown rather than an explicit shutdown call, as the existing host path already does. Replace the two author-local absolute paths in the oracle-tooling walkthrough with a user-chosen directory and a $repo variable.
c/tests/test_backend_loader.py had no CI anywhere. The python job is Linux-only, where the Windows loader contract does not apply, so 78 tests covering DLL-name selection, runtime-directory validation, physical runtime identity and the fail-closed binding decision never ran. They need no GPU and no HIP SDK: the fixtures are stub DLLs built on the fly by the MinGW gcc/objdump this job already installs. Add binutils for objdump and run them on the existing windows-latest leg. shell: pwsh overrides the job default deliberately. Under the MSYS2 shell the UCRT64 interpreter reports os.sep == "/", so the loader's backslash-joined paths compare unequal and the path tests fail; the same interpreter reports "\" from PowerShell. The step asserts that before running anything. Both MSYS2 bin directories go on PATH: ucrt64 for gcc/objdump and usr for make. Without make the host-matrix class skips at setUpClass and the run silently loses tests while still reporting OK, so the step also fails on any skip. GPU and runtime selectors are cleared so nothing can reach a real device.
The two Windows path assertions compared the loader's output against str(Path) directly. Both separator conventions can be live in one process: under MSYS2 the temp root arrives as D:/a/_temp/msys64/tmp/... with forward slashes while os.path.join adds a backslash, so the comparison fails on the separator alone. The loader is right - \ is the Windows separator and it should keep emitting it - so normalise both sides in the test instead, the same way this file already does in LoaderRuntimeBindingTest. Drop the os.sep assertion from the Windows CI step with it. It encoded the wrong diagnosis: os.sep can read \ while the temp root is still forward-slashed, so the guard passed and the tests failed anyway. The comparison no longer depends on either.
e5ff643 to
1f76e49
Compare
|
Thanks — and yes, the teacher-forced The retained authoritative sweep contained two CPU teacher-forcing processes This was an exact comparison of the selected token decisions, not a Separately, two CPU and two hybrid-HIP free-decode processes produced the same You were also right about the Windows CI failure. My earlier The fix is test-only. The loader continues to emit the correct Windows The branch is now rebased onto |
2 participants
Summary
Adds a native Windows HIP path so AMD GPUs are reachable from the Windows build:
HIP_DLL=1buildsc/coli_hip.dllfrom the existingc/backend_cuda.cuand existingcoli_cuda_*ABI. No new backend abstractionis introduced, and the Linux HIP path remains unchanged.
same seam as the existing Windows CUDA split.
COLI_HIP_RUNTIME_DIRnames the directory containingamdhip64_7.dll; theloader refuses to start the GPU tier if a different runtime is already
mapped, if more than one matching runtime is mapped, or if it cannot establish
physical file identity.
Windows CI.
gfx1151hardware.both CPU runs and both hybrid-HIP runs selected the same predicted token at
every one of the 32 checked positions.
autoregressive sequence across two independent processes per mode.
Why
HIP=1is Linux-only and is refused elsewhere by an early Makefile error. OnWindows, MinGW GCC cannot compile
.cufiles, whilehipcctargets the MSVCABI. There was therefore no native route from the Windows build to an AMD GPU;
the Windows host could select only the CUDA DLL.
The DLL split alone is insufficient because AMD ships a versioned runtime,
and Windows resolves an imported DLL basename against modules already loaded in
the process. A machine may contain several copies of
amdhip64_7.dll, such as asystem-wide ROCm installation, an SDK runtime, an unpacked development build, or
a copy in a system search location.
Silently binding to an ambient
System32or unrelated ROCm runtime is not anacceptable failure mode for a GPU backend. This implementation therefore
selects the runtime explicitly and verifies the resulting module identity.
Implementation
Native Windows HIP build
HIP_DLL=1is opt-in. The default build remains CPU-only.The HIP DLL build supports:
HIP_DLL=1HIP_ARCHHIP_SDK_ROOTHIP_BIN_DIRHIP_INCLUDE_DIRHIP_LIB_DIRHIP_DEVICE_LIB_PATHHIPCCHIP_SDK_ROOTmay default fromHIP_PATH, but every major SDK component can beoverridden independently because packaged SDK layouts do not necessarily keep
runtime, development, LLVM, and device-library files under one common root.
HIP_ARCHmust be explicit on native Windows becauserocm_agent_enumeratoris not available there.CUDA_DLL=1andHIP_DLL=1are mutually exclusive.The existing Linux HIP and Windows CUDA build paths remain supported.
Backend selection
A Windows host built with
HIP_DLL=1selectscoli_hip.dllrather thancoli_cuda.dll.The backend DLL is loaded from the host executable's directory by absolute
path. The existing
coli_cuda_*ABI remains the host/backend seam, so noparallel HIP-specific backend interface is introduced.
Deterministic runtime binding
For HIP DLL mode:
COLI_HIP_RUNTIME_DIRidentifies the directory containingamdhip64_7.dll.file index rather than path text alone.
amdhip64_7.dllis rejected.configured runtime.
initialization failure.
amdhip64_7.dllis already loadedThere is no bare-name fallback and no implicit
System32runtime search.The runtime is verified after it is mapped.
LoadLibraryExWaccepts a path,not an already-open file handle, so a file replacement between preliminary
validation and loading is detected afterward rather than prevented beforehand.
CUDA_DLL=1hosts ignoreCOLI_HIP_RUNTIME_DIRand preserve their existingcoli_cuda.dllbehaviour.Diagnostics
Windows backend-load failures now distinguish relevant failure classes,
including:
This replaces a generic failure message that could not reliably identify what
had gone wrong.
Tests
c/tests/test_backend_loader.pyis the dedicated owner for the Windowshost/backend loading contracts.
Its synthetic DLL fixtures cover:
The suite does not require ROCm, TheRock, CUDA, a physical GPU, or a real
amdhip64_7.dll.Validation
Static and regression validation
c/tests/test_backend_loader.pyunittest discover -s tests)The original 37 Python skips were environmental, including unavailable optional
Python packages, binaries not built for that test run, and POSIX-only script
coverage. None was loader-, HIP-, or backend-contract-related.
Physical Windows HIP validation
Validated configuration:
gfx1151Measured results:
C:\TheRock\build\bin\amdhip64_7.dll.System32and ROCm 7.1 copies on the same machine were not selected.gfx1151.quant_matmulkernel was executed through the productionhost/backend seam.
the CPU reference.
processes.
token-ID level.
positions in both CPU runs and both hybrid-HIP runs.
tolerance-based comparison.
20-token sequence.
positions.
qualified hybrid runs.
The model workload was a small synthetic GLM MoE:
glm_moe_dsaThe model is intended as a deterministic correctness fixture rather than a
production-model claim.
Post-rebase validation
The physical validation above was performed while the branch was based on
upstream
devat72ddb673.Upstream subsequently advanced through
44d7958to7fb1159, and this branchwas rebased onto the newer base.
The first rebase required one manual conflict resolution in
c/Makefile:tests/test_kv_prefixtarget was preservedhip-dll.PHONYentry was preservedThe resolution was the direct union of the two changes. No semantic alternative
had to be selected.
A later rebase from
44d7958to7fb1159was conflict-free.The first hosted Windows CI run also exposed two test-harness failures caused by
mixed path-separator conventions in the same process:
The loader continues to emit the correct native Windows
\separator.The fix was limited to the test and CI harness:
os.sepCI assertion was removed because it encoded an incompletediagnosis
strict about path identity rather than textual separator spelling
The resulting branch head is:
on upstream base:
Post-rebase and post-fix local validation completed with:
in each
devbaseThe physical
gfx1151model validation was not rerun on head1f76e498.No W1-authored production patch changed after the physical validation. The later
change was limited to test path normalization and the corresponding CI
diagnostic logic, while newer upstream production changes were incorporated
through the rebase.
Evidence provenance
Three harness attempts preceded the authoritative model sweep and failed
before producing qualification data:
None produced an authoritative process-result row, and their temporary logs
were not retained.
The retained qualification data consists of exactly eight final model
processes in one contiguous, correctly ordered sweep:
Every reported model comparison derives only from those eight final processes.
This is an evidence-provenance limitation, not a runtime or model failure.
CI boundary
The existing
windows-latestleg ofengines-all-platformsruns the syntheticbackend-loader contract suite.
The Windows runner may expose both separator conventions within the same
process: a temporary root may use forward slashes while the native loader
appends the runtime basename with a backslash.
The loader continues to emit the correct Windows separator. The tests normalize
both values before comparison instead of requiring identical textual spelling.
The CI step:
shell: pwshPATHusr\bintools required by the fixture buildThe skip check is important because an incomplete fixture toolchain can
otherwise cause a test class to skip while
unitteststill reports an overallsuccessful exit.
This PR does not add hosted physical Windows HIP coverage. The existing
hosted CI configuration has neither the validated Windows HIP SDK setup nor a
physical AMD GPU runner.
engine-hip-syntaxtherefore continues to cover the Linux HIP compile path,while the Radeon 8060S /
gfx1151results above were produced externally onphysical hardware and are reported as validation evidence rather than hosted CI
coverage.
Limitations
controls, including
CUDA_EXPERT_GBand an appropriate pin or usage source.That configuration was not validated here.
32 checked positions, not bit-identical logits.
control.
autoregressive decode length is 20 positions.
divergence. This is not presented as a complete isolated KV-cache proof.
explicit
coli_cuda_shutdowncall.presented as explicit backend shutdown.
successfully.
[CUDA] device 0: ...line while having no useful GPU tensorresidency.
[CUDA] resident set: N tensorswithN > 0is the useful-work signal; thedevice line alone is insufficient.
<cmath>declarations collided with HIP's
__device__overloads.VCToolsVersion=14.44.35207.will fail.
devat7fb1159.the physical model workload was not rerun on current head
1f76e498.devbase used by the current branch.
Manual operator test
No separate manual end-user/operator walkthrough was performed. The
application was nevertheless executed repeatedly through automated full-host
validation on the physical Windows machine, including model loading, teacher
forcing, and free autoregressive decoding.
Notes for reviewers
The series contains fifteen commits and keeps them separate intentionally:
deterministic runtime binding, diagnostics, synthetic fixtures, loader
contracts, and supporting tooling incrementally.
GPU_BACKENDS.mddocs/windows.md.github/workflows/ci.ymlhosted Windows run:
os.sepCI guardNo W1-authored production source changed after the physical hardware and
full-model validation.
Two PRs touch adjacent areas and may be useful context for maintainers:
c/backend_loader.c.Neither implementation was contained in the original selected base for this
branch.
Issue #759 is relevant background for integrated-GPU owners. HIP already
reports
integrated=1for the validated device, so this change does not alterthat classification.