[WIP][Fix] Fix lld invocation failed when ROCm is not at the baked-in path - #987
Open
jli-melchior wants to merge 22 commits into
Open
[WIP][Fix] Fix lld invocation failed when ROCm is not at the baked-in path#987jli-melchior wants to merge 22 commits into
lld invocation failed when ROCm is not at the baked-in path#987jli-melchior wants to merge 22 commits into
Conversation
Update build scripts to read repo/commit from internal-llvm-build-info.json and tag build/install directories with the commit short hash. Fix C++ API breakage introduced by the internal LLVM: - OpaqueProperties → PropertyRef in inferReturnTypes - Buffer op aux parameter from Value → Attribute - Atomic buffer ops now require explicit result Type - MFMA blgp parameter from IntegerAttr → MFMAPermB enum - MFMA scale cbsz/blgp from IntegerAttr → MatrixFormat enum - WMMA modC from raw uint16_t → WMMACModifier enum - WMMA scale params from raw casts → typed ROCDL enums - TDM cachePolicy from uint32_t → Attribute - DenseMapInfo: remove sentinel keys (tombstone-free hashing) Co-Authored-By: Claude <noreply@anthropic.com>
Consolidate internal-llvm-build-info.json and llvm-hash.txt into a single thirdparty/llvm-build-info.json with "upstream" and "internal" entries. Build scripts select by key name: - build_llvm.sh reads ['upstream'] - build_internal_llvm.sh / build_internal_flydsl.sh read ['internal'] Update CI workflows (build-whl, flydsl, build-custom-llvm-tools) to use llvm-build-info.json for cache keys and commit lookups. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
The internal LLVM update changed several ROCDL ops: operands became keyword-only attributes (aux, mask, cache_policy) and WMMA sign params were removed. Add compatibility wrappers in expr/rocdl so downstream callers (e.g. aiter) continue to work without source changes. - raw_ptr_buffer_load: accept aux as positional, convert Value/int → Attr - raw_ptr_buffer_load_lds: same aux handling - sched_barrier / sched_group_barrier: convert int mask → SchedGroupMask enum attr - tensor_load_to_lds / tensor_store_from_lds: accept cache_policy as positional - tdm_ops: pass cache_policy as keyword to underlying ROCDL ops Co-Authored-By: Claude <noreply@anthropic.com>
The internal LLVM update changed WMMA op parameters (fmtA, fmtB, modC, scaleAType, scaleBType, fmtScaleA, fmtScaleB) from plain integers to typed ROCDL enum attributes. Add int-to-enum conversion in the existing wmma_scale_f32_16x16x128_f8f6f4, wmma_scale_f32_32x16x128_f4, and wmma_f32_16x16x128_fp8_fp8 wrappers so downstream callers passing ints continue to work. Co-Authored-By: Claude <noreply@anthropic.com>
RawPtrBufferLoadOp/StoreOp aux changed from a positional Value operand to a keyword-only Attribute in the internal LLVM update. Convert the cache_modifier int to IntegerAttr and pass as keyword. Co-Authored-By: Claude <noreply@anthropic.com>
global_prefetch's second param changed from positional `scope` to keyword-only `cache_policy` in the LLVM binding regeneration. l2_prefetch_tile was the only call site missed — fix it to use cache_policy=_cache_policy_attr(scope), consistent with the tensor_load_to_lds/tensor_store_from_lds fixes in the same file. Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
Co-Authored-By: Claude <noreply@anthropic.com>
The upstream LLVM changed blgp from I32Attr to ROCDL_MFMAPermBAttr and cbsz/blgp in mfma_scale from I32Attr to ROCDL_MatrixFormatAttr. The AttrBuilder is not registered for these enum types, so passing a plain int causes std::bad_cast at op creation time. Add _blgp_attr() to convert int blgp values to #rocdl<mfma_perm_b ...> attributes in _split_mfma_operands, and reuse _wmma_fmt() for the mfma_scale cbsz/blgp (ROCDL_MatrixFormatAttr). Co-Authored-By: Claude <noreply@anthropic.com>
…tribute The upstream LLVM added a leading `res` (result type) parameter to raw_ptr_buffer_atomic_fadd/fmax and changed `aux` from positional Value to keyword-only Attribute. The old wrapper passed vdata as the first arg, which was misinterpreted as the result type causing std::bad_cast. Infer the result type from vdata.type, convert int aux to IntegerAttr, and pass aux as keyword. Co-Authored-By: Claude <noreply@anthropic.com>
Upstream LLVM renamed the gpu.func kernel marker attribute from "gpu.kernel" to "kernel". The _iter_gpu_kernel_funcs helper failed to match any kernel functions, so lower_compile_hints silently skipped the waves_per_eu passthrough override. Accept both "kernel" and "gpu.kernel" for backward compatibility. Co-Authored-By: Claude <noreply@anthropic.com>
Upstream LLVM now lowers rocdl.waves_per_eu directly to the LLVM IR
"amdgpu-waves-per-eu" function attribute. The old approach of deleting
the native attribute and adding a passthrough string no longer works.
Set rocdl.waves_per_eu directly with the override value. Update tests
to match the new single-value format ("2" instead of "2,2").
Co-Authored-By: Claude <noreply@anthropic.com>
Upstream LLVM/MLIR changed rocdl.mfma.scale to print cbsz and blgp as enum keywords (fp8_e4m3, fp8_e5m2, fp4_e2m1) instead of raw integers. Update the three CHECK lines in mma_atom_stateful.mlir to match.
- mma_scale_gfx1250.mlir: wmma.scale/scale16 ops now print named enum
keyword attrs (fmtA, fmtB, modC, scaleAType, ...) instead of integer
dict attrs; modC=1 prints as "neg" instead of "1 : i16".
- tdm_gfx1250.mlir: tensor.load.to.lds / tensor.store.from.lds changed
cachepolicy from keyword syntax ("cachepolicy 0") to positional (", 0").
The upstream LLVM narrowed scaleAType/scaleBType to a 2-value enum (WMMAMatrixScale: row0=0, row1=1). The old verifier accepted [0,3] but the static_cast to the enum produces malformed IR for values 2 and 3. Tighten the verifier range to [0,1] and fix the fp4_opsel FileCheck test to use opselB=1 (row1) instead of opselB=2 (out-of-range). No shipping kernels are affected — all gfx1250 WMMAScale kernels default opsel to 0. Co-Authored-By: Claude <noreply@anthropic.com>
Update test_wmma_scale_type_roundtrip to use opsel_b=1 instead of opsel_b=2, matching the tightened verifier range [0,1]. Co-Authored-By: Claude <noreply@anthropic.com>
`gpu-module-to-binary` spawns `<toolkit>/llvm/bin/ld.lld`, where `<toolkit>` comes from ROCM_PATH/ROCM_ROOT/ROCM_HOME or from `__DEFAULT_ROCM_PATH__`, a path baked into the LLVM build. Any container that installs ROCm somewhere else fails, and upstream reports only `lld invocation failed` without naming the path it tried. Our own LLVM bakes in the build machine's ROCm wheel path, so an image switch is enough to break it. Add `fly-emit-gpu-binary`, a wrapper that runs upstream only as far as `format=isa` -- which returns before the toolkit is consulted -- then finishes in process: `mlir::ROCDL::assembleIsa` for the AMDGPU MC assembler and the LLD ELF driver for the link. That removes the lookup entirely and pins the linker to the LLVM revision that produced the ISA. LLD is resolved from the same install that provides MLIR, with NO_DEFAULT_PATH so a mismatched system LLD can never be picked up, and is linked statically into the shared library. An LLVM built without the lld project keeps the upstream behavior via FLYDSL_HAS_LLD_LIBRARY. Two details are carried over from Triton's equivalent change (#7548): `--threads=1`, because LLD's use of the LLVM thread pool deadlocks in `~TaskGroup()` inside a forked child, and checking `canRunAgain`, because a JIT links thousands of times per process and must not continue on corrupted linker state. The FLYDSL_COMPILE_LLVM_DIR path still uses `gpu-module-to-binary`: it drives an upstream mlir-opt that does not know FlyDSL passes. Verified on the same IR with ROCM_PATH pointing at a nonexistent directory: upstream reports `lld invocation failed`, the new pass emits a valid ET_DYN/EM_AMDGPU `gpu.binary`. Full tests/unit + tests/system swept per file against the upstream pass: identical results. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
…ntainer Linking in process removed the ld.lld lookup, but `appendStandardLibs()` still reads `<toolkit>/amdgcn/bitcode` whenever a module references `__ocml_*` or `__ockl_*`, and errors out if that directory is absent. FlyDSL reaches it: `convert-gpu-to-rocdl` pulls in MathToROCDL, so scalar math ops with no LLVM intrinsic lower to ocml calls -- `fx.erfc` becomes `__ocml_erfc_f32`. So the per-image path problem survived in the device-library half. Bundle the bitcode the way Triton does. CMake locates ocml/ockl/hip/opencl at configure time (ROCM_PATH/ROCM_ROOT/ROCM_HOME or /opt/rocm, overridable with -DFLYDSL_ROCM_BITCODE_DIR) and copies them into the package; those four are the complete set appendStandardLibs() can request, since the oclc_* control variables are synthesized in-module rather than read from disk. They land under `_mlir/`, which is already the packaged build-output subtree, so the editable symlink and the wheel's package_dir mapping carry them with no new plumbing. `rocm_toolkit_path()` then resolves FLYDSL_COMPILE_ROCM_PATH, the bundled tree, and ROCM_PATH/ROCM_ROOT/ROCM_HOME in that order and passes the winner as `toolkit=`. Finding nothing passes no option at all, leaving upstream's lookup untouched rather than forcing a known-bad path on it. `RocmBackend.hash()` folds the bitcode digest into the JIT cache key: swapping ocml.bc changes generated code without changing any FlyDSL shared library, so the native-library hashes alone would serve stale artifacts. Verified with ROCM_PATH/ROCM_ROOT/ROCM_HOME all unset: a kernel calling fx.erfc compiles and runs. Hiding the bundled ocml.bc makes that same compilation fail on the bundled path, confirming which copy is consumed. Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
jli-melchior
changed the title
lld invocation failed when ROCm is not at the baked-in pathlld invocation failed when ROCm is not at the baked-in path
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
2 participants
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.
Motivation
Technical Details
Test Plan
Test Result
Submission Checklist