Repository navigation
[pull] master from tensorflow:master - #8853
Merged
Merged
Conversation
…ductionOp Add upper bound validation for num_segments in UnsortedSegmentReductionOp::Compute to prevent integer overflow in output tensor size calculations. Changes: - Reject num_segments > 2^31 with InvalidArgument error before allocation - Document the 2G bound rationale (safe across all architectures and dtypes) - Add Python regression test verifying num_segments > 2G raises InvalidArgumentError Fixes #117549
Addresses review feedback on #120176: the 2G guard now caps kMaxSegments at std::numeric_limits<Index>::max(), so num_segments == 2^31 is rejected for signed int32 indices instead of overflowing the static_cast<Index>. Adds an int32-index boundary regression test at exactly 2^31 and wraps an over-long assert regex to satisfy PyLint.
Addresses review feedback on #120176: - Wrap the initial read of num_segments in internal::SubtleMustCopy, as the original code did, so the value is copied out of the (possibly shared) input buffer before it is validated. Previously the copy happened after the range checks, leaving a TOCTOU window. - Reformat the changed region with clang-format (Google style) so no line exceeds 80 columns. - Restore the two blank lines before SparseSegmentReductionHelper in segment_reduction_ops_test.py.
Matches the surrounding code, which already uses absl::InvalidArgumentError throughout this file, and the style dmiltr3 had to apply by hand to get the change through the internal presubmits. The message text is unchanged, so the two regression tests still assert the same strings.
Read each num_segments scalar through internal::SubtleMustCopy inside the branch, before the widening cast, so the copy actually guards the read from the shared input buffer. Drop the redundant SubtleMustCopy on output_rows, which was applied to a local const int64_t that cannot be mutated asynchronously.
The previous bound was min(2^31, Index max), which resolved to kint32max for a 32-bit index but to 2^31 for a 64-bit one. Those differ by one, and the extra value is the unsafe one: 2^31 is not representable in a signed 32-bit integer, so it casts to -2147483648 and would wrap in any downstream code that indexes with int32. Using kint32max for all index types removes both the wrap and the asymmetry where a 64-bit index admitted one more segment than a 32-bit one, for no reason a caller could observe. Nothing practical is given up: the CPU functor allocates a std::vector<Index> row_counter(num_segments), so kint32max segments is already tens of gigabytes of bookkeeping before any output element is written. Adds testNumSegmentsOverflow_Int64Index, which pins 2^31 as rejected on the 64-bit path -- the case the old bound let through and which no test covered. Existing expectations are unaffected: testNumSegmentsOverflow (2^31 + 1), testNumSegmentsOverflow_Int32Index (2^31) and testBadNumSegments all still reject, and the error strings are unchanged.
Comment only; the bound, the error strings and the tests are unchanged. Says why kint32max is used for every index type without restating it three times, and keeps the note that the CPU functor's row_counter allocation makes the limit academic in practice.
The int32 branch now casts to int64_t rather than relying on the implicit conversion the ternary's common type performs, so both branches have the same type at the point of selection. The cast sits outside SubtleMustCopy rather than inside it. That ordering is what keeps the earlier review fix intact: SubtleMustCopy takes a const T& and reads through a volatile pointer to the address it binds, so it has to bind to the tensor buffer to be worth anything. Casting first would hand it a temporary and quietly restore the problem. Casting its result is free -- the volatile read has already happened on the buffer, and widening a copied int32 cannot lose information. No behavioural change: the ternary already yielded int64_t by promotion, so this only makes the conversion explicit and its position deliberate. Error strings are untouched and the three regression tests still assert the same messages. Also corrects a docstring that still described the bound as 2^31; it has been kint32max since 12ad001.
Add checks to ensure height and width sequence tensor values are non-negative.
…sh on dtype mismatch
SparseSegmentGradFunctor<GPUDevice, ...> (used by SparseSegmentSumGrad, SparseSegmentMeanGrad, and SparseSegmentSqrtNGrad on GPU) never validated its `indices` and `segment_ids` inputs before using them to read/write device memory. The CPU functor bounds-checks every value against the output-row count and the incoming gradient's row count before using it, returning InvalidArgumentError on failure; the GPU path had no equivalent check, so a malformed or out-of-range index caused an out-of-bounds device read/write, typically surfacing as a fatal `cudaErrorIllegalAddress` abort that also corrupts the CUDA context for the rest of the process. Add the missing validation to the GPU functor by copying the (typically small) `indices`/`segment_ids` arrays to the host and replicating the same checks and error messages the CPU functor already uses. Several existing tests for this were already restricted to `use_gpu=False` specifically to avoid the crash; this removes that restriction so they now exercise the fixed GPU path directly. Fixes #113146.
- Wrap the GPU validation logic in `if (nouter > 0)` so it no longer reads segment_vec[-1] when indices_vec is empty. - Replace the synchronous full-array copy of `indices`/`segment_ids` to the host with a device-side bounds-check kernel that scans both arrays in parallel and records only the position of a violation (if any) via GpuAtomicMin into a 2-element int32 buffer. Only that tiny buffer is transferred back to the host in the common (valid-input) case; the actual offending value is fetched with one additional single-element copy only when a violation was actually found, to build the error message. Addresses review feedback from @dmiltr3.
The bounds-check kernel packs each violation's position into an int32 buffer via GpuAtomicMin. If Index is int64_t and nouter exceeded INT32_MAX, a position beyond INT32_MAX would truncate to a negative value when narrowed, which would then be used as a negative pointer offset when fetching the offending value for the error message. Guard against this the same way SparseSegmentGradV2Functor already does for its own nouter/ninner (this file, above): reject the input with a clear error instead of proceeding with a value that doesn't fit.
…ive memory space color. PiperOrigin-RevId: 987355696
Imported from GitHub PR openxla/xla#49314 Bumps [astral-sh/setup-uv](https://github.com/astral-sh/setup-uv) from 10.0.1 to 10.1.0. <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/astral-sh/setup-uv/releases">astral-sh/setup-uv's releases</a>.</em></p> <blockquote> <h2>v10.1.0 🌈 New output <code>python-runtime-id</code>and respect NO_PROXY</h2> <h2>Changes</h2> <p>This release adds more bheind the scene security improvements and also 2 small improvements.</p> <h3>NO_PROXY</h3> <p>This action now respects <code>no_proxy/NO_PROXY</code> environment variables which were previously ignored.</p> <h3>New output <code>python-runtime-id</code></h3> <p>The new output <code>python-runtime-id</code> can be used to know which python version exactly was installed if you use <code>activate-environment</code>. See <a href="https://redirect.github.com/pyca/cryptography/pull/15572#discussion_r3913508686">pyca/cryptography#15572</a> for details on why this can be useful.</p> <h2>🐛 Bug fixes</h2> <ul> <li>fix: respect no proxy directive <a href="https://github.com/mj0nez"><code>@mj0nez</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1037">#1037</a>)</li> <li>Use JSON + a typed wrapper instead of TS codegen <a href="https://github.com/woodruffw"><code>@woodruffw</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1025">#1025</a>)</li> </ul> <h2>🚀 Enhancements</h2> <ul> <li>Expose a Python "identity" output <a href="https://github.com/woodruffw"><code>@woodruffw</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1036">#1036</a>)</li> <li>Verify downloads with astral-sh/versions checksums <a href="https://github.com/zaniebot"><code>@zaniebot</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1033">#1033</a>)</li> </ul> <h2>🧰 Maintenance</h2> <ul> <li>chore: update known checksums for 0.12.12 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1041">#1041</a>)</li> <li>chore: update known checksums for 0.12.10/0.12.11 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1038">#1038</a>)</li> <li>chore: update known checksums for 0.12.9 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1035">#1035</a>)</li> <li>chore: update known checksums for 0.12.7/0.12.8 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1031">#1031</a>)</li> <li>chore: update known checksums for 0.12.6 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1030">#1030</a>)</li> <li>chore: update known checksums for 0.12.5 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1020">#1020</a>)</li> <li>Use self-repo syntax for all in-repo actions/reusable workflows <a href="https://github.com/woodruffw"><code>@woodruffw</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1024">#1024</a>)</li> <li>Pin one-shot tools <a href="https://github.com/woodruffw"><code>@woodruffw</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1022">#1022</a>)</li> <li>ci: remove obsolete direct push attempts <a href="https://github.com/eifinger"><code>@eifinger</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1019">#1019</a>)</li> </ul> <h2>📚 Documentation</h2> <ul> <li>docs: update version references to v10.0.1 @<a href="https://github.com/apps/github-actions">github-actions[bot]</a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1018">#1018</a>)</li> </ul> <h2>⬆️ Dependency updates</h2> <ul> <li>chore(deps-dev): roll up Dependabot updates <a href="https://github.com/eifinger"><code>@eifinger</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1043">#1043</a>)</li> <li>Harden npm install defaults <a href="https://github.com/zaniebot"><code>@zaniebot</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1026">#1026</a>)</li> <li>Add dependency cooldowns <a href="https://github.com/woodruffw"><code>@woodruffw</code></a> (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1021">#1021</a>)</li> </ul> </blockquote> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/astral-sh/setup-uv/commit/bec219d24cd3e171d82865faccec33120bb574f4"><code>bec219d</code></a> chore(deps-dev): roll up Dependabot updates (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1043">#1043</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/b90ec40d15bfa44c33c6700196eb6efcdddb4373"><code>b90ec40</code></a> fix: respect no proxy directive (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1037">#1037</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/421feb646df5262e7dd93bc54161edfa30372417"><code>421feb6</code></a> chore: update known checksums for 0.12.12 (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1041">#1041</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/f634bf473ad85bf3e23a613f52c5fa9f363874fc"><code>f634bf4</code></a> Expose a Python "identity" output (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1036">#1036</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/a6772c8f0a09dc9e3582c70a994b0c55af921803"><code>a6772c8</code></a> chore: update known checksums for 0.12.10/0.12.11 (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1038">#1038</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/e105c8fb1d7b13074b851babdaef4185243c6a07"><code>e105c8f</code></a> chore: update known checksums for 0.12.9 (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1035">#1035</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/cd13f9217092d43a771cf9ba7b09bdd3da8d7c4d"><code>cd13f92</code></a> Verify downloads with astral-sh/versions checksums (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1033">#1033</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/3aef7b92c52cec135792ea1e95f4c77683d39e61"><code>3aef7b9</code></a> chore: update known checksums for 0.12.7/0.12.8 (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1031">#1031</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/d08d816a1ea176d61a318eff45abd3dffef415b1"><code>d08d816</code></a> chore: update known checksums for 0.12.6 (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1030">#1030</a>)</li> <li><a href="https://github.com/astral-sh/setup-uv/commit/19b4d1e990bec64818914c40230bde93a0de300b"><code>19b4d1e</code></a> Harden npm install defaults (<a href="https://redirect.github.com/astral-sh/setup-uv/issues/1026">#1026</a>)</li> <li>Additional commits viewable in <a href="https://github.com/astral-sh/setup-uv/compare/20cfd1bf945f4377ade1205e4dbc17946fc9a30d...bec219d24cd3e171d82865faccec33120bb574f4">compare view</a></li> </ul> </details> <br /> [](https://docs.github.com/en/github/managing-security-vulnerabilities/about-dependabot-security-updates#about-compatibility-scores) Dependabot will resolve any conflicts with this PR as long as you don't alter it yourself. You can also trigger a rebase manually by commenting `@dependabot rebase`. [//]: # (dependabot-automerge-start) [//]: # (dependabot-automerge-end) --- <details> <summary>Dependabot commands and options</summary> <br /> You can trigger Dependabot actions by commenting on this PR: - `@dependabot rebase` will rebase this PR - `@dependabot recreate` will recreate this PR, overwriting any edits that have been made to it - `@dependabot show <dependency name> ignore conditions` will show all of the ignore conditions of the specified dependency - `@dependabot ignore this major version` will close this PR and stop Dependabot creating any more for this major version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this minor version` will close this PR and stop Dependabot creating any more for this minor version (unless you reopen the PR or upgrade to it yourself) - `@dependabot ignore this dependency` will close this PR and stop Dependabot creating any more for this dependency (unless you reopen the PR or upgrade to it yourself) </details> Copybara import of the project: -- 454adc369ab7b2a599e818b54fdb1ced4f464ea6 by dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>: Bump astral-sh/setup-uv from 10.0.1 to 10.1.0 Bumps [astral-sh/setup-uv](https://github.com/astral-sh/setup-uv) from 10.0.1 to 10.1.0. - [Release notes](https://github.com/astral-sh/setup-uv/releases) - [Commits](astral-sh/setup-uv@20cfd1b...bec219d) --- updated-dependencies: - dependency-name: astral-sh/setup-uv dependency-version: 10.1.0 dependency-type: direct:production update-type: version-update:semver-minor ... Signed-off-by: dependabot[bot] <support@github.com> Merging this change closes #49314 PiperOrigin-RevId: 987366858
Imported from GitHub PR openxla/xla#49372 📝 Summary of Changes - Enable `//xla/backends/gpu/tests:vectorization.hlo.test` on mi200 and mi350. Same `gpu/` lit as CUDA. Removes those two specs from the disabled list. - FileCheck LLVM `store <4 x i8>` at `--stage=llvm-before-optimizations`, the same stage and width as the H100/A100 lines. - Gate `--stage=ptx` with `%if !IS_ROCM`. `AMDGPUCompiler::CompileTargetBinary` returns an empty asm text, and `hlo-opt` has no GCN stage. NVIDIA PTX checks stay. 🎯 Justification Explain why this change is important and which workload benefits from this change. 🚀 Kind of Contribution 🧪 Tests Copybara import of the project: -- 9b122c87b8c8923be80148b641034d1a5330d3ad by linchen1 <lin.chen1@amd.com>: Enable vectorization.hlo check on ROCm mi200 and mi350. Signed-off-by: linchen1 <lin.chen1@amd.com> Merging this change closes #49372 PiperOrigin-RevId: 987373576
PiperOrigin-RevId: 987373870
…ure-mismatch PiperOrigin-RevId: 987377724
…ecutors when no supported devices are found. This simplifies debugging for example when VMM API is not available. See: openxla/xla#49252 PiperOrigin-RevId: 987377773
…ang-offload-bundler target IDs. Imported from GitHub PR openxla/xla#49033 📝 Summary of Changes - Add rocm-only //xla/stream_executor/gpu:gpu_hsaco_bundle_test: compile inlined AMDGPU LLVM IR with amdgpu::CompileToHsaco, then pack the HSACO with BundleGpuAsm(HsacoImage). - Fix ROCm 7 clang-offload-bundler --targets= IDs in gpu/asm_compiler.cc (4-field host and HIP triples). - Does not tag-drop //xla/stream_executor/cuda:subprocess_compilation_no_fakes_test (cuda-only). This is extra coverage, not a name-matched twin. 🎯 Justification Add more unit test coverage on ROCm platform. 🚀 Kind of Contribution 🧪 Tests Copybara import of the project: -- 981e93ac52fbeef919ccda3fdaf17cd7a2c6dd8f by Lin Chen1 <lin.chen1@amd.com>: Add ROCm HSACO compile-then-pack coverage and fix clang-offload-bundler target IDs. Signed-off-by: Lin Chen1 <lin.chen1@amd.com> Merging this change closes #49033 PiperOrigin-RevId: 987378077
PiperOrigin-RevId: 987380727
…target visibility Replace `"//xla/backends/gpu/libraries/native_custom_call_thunks:handler_allowlist"` in the `visibility` lists of core non-GPU XLA targets (`hlo`, `compiler`, `buffer_assignment`, `kernel_arguments`, `launch_dim`, `kernel_spec`) with `"//xla:native_custom_call_handler_allowlist"` and include `native_custom_call_thunks/...` in `native_custom_call_handler_allowlist` so non-GPU XLA packages do not reference a `package_group` inside `xla/backends/gpu/`. PiperOrigin-RevId: 987404559
…-zero-size-oob PiperOrigin-RevId: 987415547
CPU targets without specific ISA extensions (such as AVX512-BF16) lack legal vector register classes for narrow floats (e.g., bf16, fp16). Using these types causes LLVM to scalarize them and generate libcalls. PiperOrigin-RevId: 987420616
Add pre-checks for P2P symmetric memory support via cudaDeviceCanAccessPeer, post-checks for null windows returned from ncclCommWindowRegister, and defensive null window checks in multimem_addr, peer_addr, and the destructor. PiperOrigin-RevId: 987435182
Adds cudnn_scaled_dot_test covering block-scaled dot support in cuDNN on Blackwell. Parameterized across MXFP8 and NVFP4 type combinations and transposition layouts, reporting lowering stages (composite, optimized HLO, cuDNN graph representation) and execution parity against reference decomposition. PiperOrigin-RevId: 987447766
PiperOrigin-RevId: 987452514
AlgebraicSimplifier can rewrite bitcast-convert instructions into bitcast even when the element bitwidth changes (when the minor-most dimension is contiguous). While PriorityFusion and GpuFusible guard against fusing width-changing bitcasts using hlo_instruction_utils::KeepsBitwidth, CopyFusion previously skipped through any single-user bitcast without checking its bitwidth. Fusing a width-changing bitcast and its consuming copy into a producer fusion breaks the MLIR fusion emitters because GetBitcastMap requires matching element counts between roots and elemental lowering cannot emit a scalar mhlo.bitcast_convert across different bitwidths. Guard kBitcast checks in CopyFusion with hlo_instruction_utils::KeepsBitwidth. PiperOrigin-RevId: 987463670
Now that we have this compatibility window, we'll often need to ask OSS developers to change their PRs to avoid breaking it (e.g. openxla/xla#46865). So we should have some documentation that we can point to. PiperOrigin-RevId: 987466715
…al bugs PiperOrigin-RevId: 987488947
That will be used later to add constraints to one indexing map. Should also marginally imporove compilation time. Also call simplify in tiling propagation only if we are not in "symbolic" case where we will not use them simplified anyway. PiperOrigin-RevId: 987492063
…ing. `HandleProtoMergeRecord` read each record into a flat `absl::string_view` and then called `MergeFromString`. For records larger than a riegeli buffer (e.g. an HLO module carrying hundreds of MB of constants) the decoded record spans several blocks, and producing a flat view forces riegeli to copy it into a contiguous scratch buffer first. Read the record as a `riegeli::Chain` instead, which shares the decoded blocks, and parse it with `riegeli::ParseMessage(..., set_merge(true))`, which consumes the chain directly. PiperOrigin-RevId: 987496597
`serialized_executable` is already a snappy-compressed split `GpuExecutableProto` (see `WriteSplitGpuExecutable`). Compressing the outer container again saves no space (dropping it adds 0.06% on a 330 MB executable), but it costs an extra compression pass on write and a full decompression pass on load. Riegeli stores the compression type per chunk, so readers don't change and existing artifacts still load. Alternative considered: leave the inner proto uncompressed and compress the outer one instead. That would mean deciding per use case when to skip compression, which adds logic for little gain, since the outer proto is small apart from `serialized_executable`. PiperOrigin-RevId: 987497408
`Future<std::vector<T>>({})` cannot select the templated converting
constructors of `Future`, because a braced-init-list cannot deduce `U`.
Overload resolution therefore falls back to the move constructor from a
value-initialized `Future`, and both typed `JoinFutures` overloads returned an
invalid future for an empty span instead of a ready empty vector. Calling
`Map`, `OnReady` or `Await` on it dereferences a null `AsyncValue`.
PiperOrigin-RevId: 987497834
Provides a script to automate regenerating and promoting AOT compatibility golden snapshots from Bazel undeclared test outputs into `<package>/executables/<target_name>/v<N+1>/`. Also updates `AOTInterceptionPjrtClient::VerifyAgainstGolden` and `test_lib.cc` (`GetUndeclaredOutputsDir`) to reference `update_goldens.py` in golden comparison failure and precondition error messages. The golden comparison failure message prints the failing test's full label (from `TEST_TARGET`), which is what the script takes as its argument. PiperOrigin-RevId: 9875225
Goldens move to `executables/collective_ops_aot_test_2gpu_<arch>/`, one directory per backend. gb300 stays excluded through the existing `disabled_backends` entry (b/491194726); no gb300 target is generated, so `DetectGpuArchToken()` has no gb300 branch and a gb300 device would hit its `LOG(FATAL)`. PiperOrigin-RevId: 987552652
PiperOrigin-RevId: 987558515
…hile loop transformations - In WhileUtil::AppendToWhileLoopOriginalValue and WhileUtil::MakeInstructionsLiveIn, update while instruction, while init, body root, body parameter, and condition parameter to keep all loop instruction OriginalValues structurally consistent when widening while loops, and preserve call_hierarchy. - In WhileLoopSimplifier (TryRemoveDeadTupleElements, TryRemoveConstantTupleElements, TryFlattenNestedTuples, and TrySimplifyInductionVariables), propagate and structurally update OriginalValues across all 5 loop boundary instructions and verify compatibility with the new shapes. PiperOrigin-RevId: 987569790
When while loops are unrolled the same thunk is called with `Record` multiple times for different loop iterations. Currently `CustomCallRecordState` stored a single sequence of `XLA_FFI_Command*`, so during update only the last recorded sequence of commands was passed to the FFI record API, overwriting the last iteration and leaving earlier iterations of the while loop with stale buffer pointers. Keying the state by the sink command pointer returned on `RecordCreate` (and passed back via `RecordUpdate::command`) so each unrolled iteration updates its own recorded command sequence. PiperOrigin-RevId: 987596254
…verflow PiperOrigin-RevId: 987596387
Reverts 2a35254 PiperOrigin-RevId: 987600180
Also includes a quant_spec metadata in tfl.fully_connected to allow encoding custom specs. PiperOrigin-RevId: 987624321
PiperOrigin-RevId: 987625577
…on, and transfers. - Introduce dedicated bounded thread pools in `PjRtCpuRawClient` for compilation (`compile_thread_pool_`) and execution (`execute_work_runner_`), separate from the general `async_work_runner_` used for buffer transfers and linearization/delinearization. - Route `CompileInternal` to `compile_thread_pool_` and `CpuPjRtRawLoadedExecutable::Execute` to `execute_work_runner_`. Separating execution, compilation, and transfer thread pools prevents multi-device collective launches from starving antecedent input buffer transfers on `async_work_runner_`, while keeping all three pools bounded so bursts of D2H/H2D buffer transfers remain throttled. PiperOrigin-RevId: 987628676
…latforms.
Introduce `HloOpcode::kShuffle` (`HloShuffleInstruction`), which moves the
elements of its operand around along a set of `dimensions`. A single opcode
covers the whole family of data-movement patterns (rotate, reverse, permute,
...), which keeps the semantic intent available to compiler analysis, layout
assignment, and target-specific lowerings instead of forcing an expansion into
slice and concatenate during graph construction.
The pattern to apply is selected by `xla.ShuffleMode`, a `oneof` pairing each
mode with exactly the attributes that parameterize it, so a mode cannot be
combined with another mode's attributes and adding a mode adds no fields to
`HloInstructionProto`. Switches over `mode()` are exhaustive and have no
`default`, so a new mode does not compile until every consumer has handled it.
Only `rotate` is implemented here. It takes one `shifts` entry per shuffled
dimension and rotates the elements to the left.
Specifically:
* **HLO IR & Parser**: `HloShuffleInstruction` carries `dimensions` and a
`ShuffleMode`, with text serialization and parsing. The text form is
`shuffle(x), dimensions={0,1}, mode=rotate, shifts={2,5}`, where the name of
a mode is the name of its field in `ShuffleMode`.
* **Shape Inference & Verifier**: `ShapeInference::InferShuffleShape` and the
verifier reject duplicate or out-of-range `dimensions`, a missing mode, and
mode attributes inconsistent with the mode (for `rotate`, `dimensions` and
`shifts` of unequal size).
* **XlaBuilder API**: `xla::Shuffle` constructs shuffles from symbolic
handles. The new `shuffle` utility provides `MakeRotateMode` to build a rotate
mode and `NormalizeShift` to map an arbitrary shift onto `[0, dim_size)`.
* **HLO-to-MHLO / StableHLO Translation**: `HloFunctionImporter` imports
`kShuffle`, lowering each rotated dimension sequentially into
`stablehlo.slice` and `stablehlo.concatenate`
(`concat(slice(shift:), slice(0:shift))`).
* **Pass Integration**: Register `kShuffle` across core compiler passes
including instruction fusion, layout assignment, and sharding propagation.
PiperOrigin-RevId: 987630885
PiperOrigin-RevId: 987639736
…core/tfrt/ifrt:ifrt_device_utils_test PiperOrigin-RevId: 987641321
PiperOrigin-RevId: 987716993
Extend the HLO optimization pipeline with algebraic simplifications for `HloOpcode::kShuffle` in `kRotate` mode, allowing compiler pipelines to simplify rotations. Specifically, this commit introduces: * **AlgebraicSimplifier (`algebraic_simplifier.cc`)**: Add simplification rules for `kShuffle::kRotate` instructions - Eliminating identity rotations (`norm_shift == 0` , empty shift dimensions, splat dimension) - Folding/combining chained `kShuffle::kRotate` operations. - Canonicalize shifts to be in `[0, dim_size)` and order dimensions in ascending order * **Testing**: Add comprehensive unit tests in `AlgebraicSimplifierTest`. PiperOrigin-RevId: 987747347
…rtitioning. After GSPMD generates device local, it clears sharding except for two cases: (1) Send/Recv instructions and (2) Parameter instructions. Case (1) is due to the fact that GSPMD doesn't handle the single device sharding on Send/Recv instructions, and it leaves the sharding for later passses to handle it. Case (2) is due to the fact that there are many places in the compilation pipeline that find the module input sharding from the sharding on parameter instructions, even though the information is also available in the module spmd_parameters_shardings metadata. This CL makes GSPMD also clear sharding on parameter instructions, and updates downstream users to read parameter shardings from the module's spmd_parameters_shardings metadata. PiperOrigin-RevId: 987756729
…ySpaceAssignment
Root Cause:
`MemorySpaceAssignment::Process()` rewrites the HLO graph in two sequential loops over all `Allocation` objects:
1. Loop 1 calls the virtual method `Allocation::Process()` (`CopyAllocation::Process()`, `PinnedAllocation::Process()`, and `ParentAllocation::Process()`).
2. Loop 2 calls the virtual method `Allocation::PostProcess()` (`ParentAllocation::PostProcess()`).
Several HLO transformations across these two loops failed to maintain or propagate `OriginalValue`:
1. While-loop widening in `ParentAllocation::Process()` (Loop 1) and `ParentAllocation::PostProcess()` (Loop 2):
When `ParentAllocation` passes a buffer from a parent computation into a `while` loop, it widens the `while` loop tuple from N elements to N+1 elements at `new_tuple_index`:
- `ParentAllocation::Process()` widens the `while` loop inputs by creating a new input tuple (`new_while_operand = TupleUtil::ReplaceTupleWith(...)`) and mutating the shapes of `calling_instruction_` (the `kWhile` instruction), `while_condition->parameter_instruction(0)`, and `while_body->parameter_instruction(0)` in-place. Because the in-place mutated instructions kept their old N-element `OriginalValue` trees while their shapes became (N+1)-tuples, `HloVerifier` failed with shape mismatches. Meanwhile, `new_while_operand` was created with `original_value() == nullptr`.
- To keep downstream users of the `while` loop valid, `ParentAllocation::Process()` creates `tuple_with_old_shape = TupleUtil::ExtractPrefix(calling_instruction_, new_tuple_index)` and replaces uses of `calling_instruction_` with `tuple_with_old_shape`. `tuple_with_old_shape` and its `GetTupleElement` operands were created with `original_value() == nullptr`, dropping value tracking for all users after the loop.
- `ParentAllocation::PostProcess()` runs in Loop 2 (after all Loop 1 allocations inside `while_body` have finished) to widen the `while_body` root tuple (`new_while_body_root = TupleUtil::ReplaceTupleWith(added_element, old_body_root, ...)`), which also left `new_while_body_root` with `original_value() == nullptr` and dropped the appended element's `OriginalValue` at `{new_tuple_index}`.
2. Tuple replacement and async copy creation in `CopyAllocation::Process()` and `Allocation::UpdateUses()` (Loop 1):
- `CopyAllocation::Process()` inserts `copy-start` and `copy-done` instructions to move buffers between default memory (HBM) and alternate memory (VMEM), but left `copy_done_` with `original_value() == nullptr` instead of copying `producing_instruction`'s `OriginalValue`.
- `CopyAllocation::Process()`, `PinnedAllocation::Process()`, and `ParentAllocation::Process()` all call the helper `Allocation::UpdateUses()` during Loop 1 to rewire consumer instructions to use `producing_instruction`. When a consumer is a `kTuple` instruction (such as `while_body->root_instruction()` or `while_init`), `Allocation::UpdateUses()` calls `TupleUtil::ReplaceTupleWith`, which constructs a brand-new `kTuple` instruction with `original_value() == nullptr`. Consequently, any Loop 1 allocation updating an element of `while_body->root_instruction()` wiped out its `OriginalValue` to `nullptr` before Loop 2 (`ParentAllocation::PostProcess()`) even ran.
What This CL Fixes:
- Add `TupleTree::num_children(ShapeIndexView index = {})` in `tuple_tree.h` to query the number of children of a tuple node directly from the index table.
- Introduce `AppendToTupleOriginalValue(src_tuple, dest_tuple, appended_instr)` in `allocation.cc`, which uses `CopyOriginalValue` from `hlo_original_value_util.h` to map existing tuple elements from `src_tuple` to the widened `dest_tuple` and copies `appended_instr->original_value()->tree()` into the last element (`dest_tuple->shape().tuple_shapes().size() - 1`) via `CopyCompatibleSubtreeFrom`.
- In `ParentAllocation::Process()`, before widening `calling_instruction_` in-place, copy its N-element `OriginalValue` to `tuple_with_old_shape` via `CopyOriginalValue` and propagate subtrees to its `GetTupleElement` operands. Then widen `new_while_operand` with `AppendToTupleOriginalValue` and update `calling_instruction_`, `while_condition` parameter 0, and `while_body` parameter 0 via `AppendToWhileLoopOriginalValue`.
- In `AppendToWhileLoopOriginalValue` (`while_util.cc`), only update `while_body->root_instruction()` if its shape has already been widened to match `while_instr->shape()`, supporting two-phase while-loop widening where `while_body` root widening is deferred to `ParentAllocation::PostProcess()`.
- In `Allocation::AddGetTupleElements()`, propagate the `OriginalValue` subtree from `defining_position().instruction` at `defining_position().index` to the newly created `GetTupleElement` instruction.
- In `ParentAllocation::PostProcess()`, call `AppendToTupleOriginalValue(old_body_root, new_while_body_root, added_element)` so the widened `while_body` root tuple preserves both the body root's `OriginalValue` and the pass-through element's `OriginalValue`.
- In `CopyAllocation::Process()`, propagate `producing_instruction`'s `OriginalValue` to `copy_done_` (and any intermediate bitcast).
- In `Allocation::UpdateUses()`, when `TupleUtil::ReplaceTupleWith` replaces a `kTuple` instruction, clone `tuple_inst`'s `OriginalValue` onto `replacement_instruction` and update `{use.operand_index}` from `producing_instruction->original_value()`.
PiperOrigin-RevId: 987774950
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 subscribe to this conversation on GitHub.
Already have an account?
Sign in.
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.
See Commits and Changes for more details.
Created by
pull[bot] (v2.0.0-alpha.4)
Can you help keep this open source service alive? 💖 Please sponsor : )