Repository navigation
Conversation
|
Thanks for opening a PR! This project follows LLVM's incremental-development and AI-tool-use Before requesting review, please check that:
Reviewers are assigned through |
workflow_dispatch requires the workflow to exist on the default branch, which leaves the selfbuild job unreachable from a fork topic branch. pull_request makes the PR itself run it, and FROM_STEP now carries a default because that event supplies no inputs. windows-build and linux-build are gated to ROCm/hip-ep so a fork PR does not start llvm/ort/oga or the StrixHalo gpu_test. Also drop zlib's example/test targets from the third-party step.
The runner's powershell shell sets ErrorActionPreference=Stop, which turns a redirected native stderr into a terminating NativeCommandError. rev-parse HEAD on a freshly initialized repo always writes there, so the clone step died before any fetch ran. Use --verify --quiet and gate on $LASTEXITCODE, and keep stderr non-fatal inside the helpers so curl progress, cmake warnings and git fetch output cannot abort a step either.
zlib names its static archive zs.lib, while CMake's FindZLIB only searches zlibstatic/zlibstat/zlib/z, so LLVM failed configure with "Could NOT find ZLIB (missing: ZLIB_LIBRARY)" even though the header resolved. Copy the archive to zlibstatic.lib after install. Key the third-party step on that archive instead of zlib.h so a prefix that only has the header is not treated as provisioned.
The runner's SCCACHE_DIR is a NAS share, which would make every TU lookup an SMB round trip across an LLVM build. Start a dedicated server on a local disk under SCCACHE_SERVER_PORT so the shared server other jobs rely on is left alone, and assert the reported cache location actually moved because a running server keeps its original directory regardless of the environment.
sccache reports its cache path with escaped separators, so the assert rejected a server that had in fact moved to the local disk. The sdk-selfbuild tree stays persistent so from_step can resume, which also means nothing ever reclaims it. clean=build drops the build trees and the SDK prefix while keeping the clones; clean=all drops the tree outright.
Resumable builds traded disk and stale state for wall time that sccache already provides, so drop from_step and the clean input: build trees, the SDK prefix and the third-party prefix are removed at the start of every run. Pinned clones survive since they are re-fetched only when a pin moves. Also stop relocating the cache. sccache now runs with whatever directory the runner is configured with, on the shared server, as it did before.
hipamd's install does a REMOVE_RECURSE on the prefix include directory on Windows, so amd_comgr.h disappeared between the comgr and hipBLASLt steps and stinkytofu's ComgrProbe.cpp failed to compile. The lib/cmake export survives that wipe, which is why find_package still succeeded and the failure only surfaced during compilation. Re-run the comgr install after hip-clr and assert the header is back. Also check it in the verify step so a future regression is caught before a build gets that far.
2 of 8 tasks
minwindef.h defines max(a,b), so numeric_limits::max() in the HIP fp8 header and throughout Tensile's headers expanded as a macro invocation with too few arguments and rocblaslt_auxiliary.cpp failed to compile. hipBLASLt's own tasks.py adds both flags for Windows builds, the second one because amdclang++ does not read ROCM_PATH there, so mirror that pair rather than only silencing the macro.
stinkytofu builds its tablegen host tools with -Werror, where clang turns the unused --rocm-path into an error on every non-HIP translation unit. hipBLASLt's tasks.py gets away with the flag because that path does not enable -Werror. Nothing has shown the flag to be necessary: the five earlier steps build without it, and the flat layout lets clang derive the ROCm root from its own location. NOMINMAX, which the previous commit was actually about, stays.
The trim only filtered bin\*.exe, so a full LLVM install left libclang.dll, LTO.dll and Remarks.dll behind: 157 MB of the 742 MB artifact that nothing consumes. Filter bin\*.dll against the same kind of keep-list, matching by pattern because the runtime DLLs carry ROCm version suffixes, and add hipblaslt.dll to the post-trim presence check. Also report sccache stats when entering the hipBLASLt step. Tensile drives clang directly for most of that step, so the server reaches its idle timeout and the final stats read zero even though the cache was hit throughout.
The artifact was a directory, so GitHub wrapped it in its own zip and a download produced loose files that still had to be repacked before they could go on a release. cmake/deps.cmake fetches one .tar.gz and unpacks it with --strip-components=1, so produce exactly that and upload the archive itself, with artifact compression off since the payload is already gzipped.
The SDK shipped an unsuffixed amd_comgr.dll next to hiprtc0716.dll, unlike every released ROCm layout. Derive the suffix from projects/hip/VERSION, the same file clr reads for hiprtc, so the two stay in step without hardcoding digits that move with the pin. rocclr bakes COMGR_DLL_NAME in and LoadLibrary's it at runtime, so hip-clr has to be told the same name; otherwise amdhip64 would look for the default while the package only carries the suffixed one, and nothing would fail until a session tried to start.
The artifact carried the LLVM pin and no ROCm version, so it read nothing like the entries cmake/deps.txt consumes. Build the name as <what>-dist-windows-<archs>-<version>-<variant> instead, taking the version from the HIP source tree so it tracks the pin. Naming now happens in the pack step rather than up front, because the HIP VERSION file only exists after the clones. The arch tag and the variant are workflow env values: a fork that changes the target set or the hipBLASLt trimming edits one line each.
The concurrency group spans branches on purpose, since every run shares one prefix and build tree under runner.workspace. With cancel-in-progress off, a group still keeps only one pending run, so pushing twice in a row left a trail of runs cancelled while queued rather than one build of the newest commit. Cancelling the in-progress build costs little now that sccache carries a 94% hit rate across runs.
A group spanning branches meant the two stacked PRs fought over the single pending slot, so pushing to one discarded the other's queued build. Builds do not actually need the shared group to serialize: strix_build is a single self-hosted runner, so the second job waits in the runner queue, where it is not at risk of being cancelled. Pushing again to the same branch still supersedes its own run. Note that the shared prefix and build tree under runner.workspace now rely on that single runner, so a second one with this label would need per-branch directories.
The Tensile .co are compressed clang offload bundles, and the SDK's own clang-offload-bundler picks the format: with LLVM_ENABLE_ZSTD=OFF it silently falls back to zlib. Against the pinned TheRock tarball that costs 3x on the gfx1151 keep-set stems the two packages share, 1720 KB against 5134 KB, while the bundles decompress to within 2.2% of each other. gfx1170 pays the most at 25.8 MB for four stems. Build zstd 1.5.7 static alongside zlib and configure LLVM with FORCE_ON, so a missing zstd fails configure instead of shipping zlib bundles. The verify step now reads the CCOB compression method out of every Tensile object, the check the earlier fallback slipped past. comgr picks zstd up through LLVMConfig's find_dependency, which is what makes this safe to ship: comgr is what unbundles a .co at hipModuleLoad, so zstd objects would not load against a comgr built without it.
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Summary
Adds a fork-only Windows workflow that builds the hip-ep ROCm SDK from source with six standalone cmake invocations and uploads a flat-layout artifact, without the TheRock superbuild.
Related issue or design
None. Fork-local counterpart of ROCm#917; this PR stays on
zz002/hip-epand is not for upstream.Why
Producing the trimmed Windows SDK through TheRock also builds components the release whitelist discards (rocBLAS, rocRAND, composable_kernel, host-blas) and relies on its flatten step, while most of the wall time sits in LLVM with no way to resume a failed step. Building only the six components we ship, directly into one flat prefix, removes both the flatten and the trim stage.
The flat layout also makes the two hipcc patches unnecessary. Both are keyed on
lib/llvm/bin, which a flat prefix does not have:0002would pointconstructCompilerPathat a directory that does not exist, and0006's guard never fires so it falls through toHIP_PATHexactly like the unpatched code. Unpatched hipcc resolvesgetHipPath()/bin, which is where this workflow installs clang.What
windows-sdk-selfbuild.yml: pinned clones, third-party zlib / zstd / msgpack-cxx / rocm-cmake,dvc pullfor the prebuilt PAL and wkmi libs, then llvm, hipcc, AMDDeviceLibs, amd-comgr, hip-clr and hipBLASLt as six separate cmake configure/build/install rounds intorunner.workspace. Source pins are workflowenvvalues (llvm-project2904a87d, rocm-systemsb0bc2cc5, rocm-librariesc37db3c2, rocm-cmake10155d72). Targets aregfx1150;gfx1151;gfx1152;gfx1153;gfx1170.HIPBLASLT_LIBLOGIC_PATHplusTENSILELITE_LOGIC_FILTER=*); staging the YAMLs without the filter still builds the full set. Generator parallelism is capped throughTENSILELITE_BUILD_PARALLEL_LEVELrather than by patching Tensile.FORCE_ON. The Tensile.coare compressed clang offload bundles and the bundler silently falls back to zlib without zstd: on the gfx1151 keep-set stems this package shares with the pinned TheRock tarball that is 5134 KB against 1720 KB, for bundles that decompress to within 2.2% of each other. The verify step now reads the CCOB compression method out of every Tensile object. comgr inherits zstd through LLVMConfig'sfind_dependency, which it must, since comgr is what unbundles a.coathipModuleLoad.windows-build.ymlandlinux-build.ymlongithub.repository == 'ROCm/hip-ep'so this fork PR does not start llvm/ort/oga or theStrixHalogpu_test.Test plan
Windows SDK Selfbuildreachesstrix_buildand all six cmake steps succeed on a cold prefixbin/clang.exe,bin/amdgpu-arch.exe,bin/hipcc*,amdgcn/bitcode,lib/clang,bin/amdhip64*.dll,bin/hiprtc*.dll,bin/amd_comgr*.dll,bin/hipblaslt.dllandlib/cmake/hip/hip-config.cmakebin/hipblaslt/library/gfx*directory holds at most 20.cofiles,gfx1170includedNOT ZSTDobject, and the gfx1151.cototal drops from 5.8 MB to roughly 1.8 MBNotes for reviewers
This round omits rocBLAS, rocRAND, composable_kernel and SPIRV and applies no patches. Tensile logic for gfx1170 exists only in
AMD-ROCm-Internal/rocm-libraries, so the private clones still depend on the runner's ownghlogin rather than the job token.Checklist