Skip to content

[pull] master from tensorflow:master - #8853

Merged
pull[bot] merged 53 commits into
Cache-Cloud:masterfrom
tensorflow:master
Sep 26, 2026
Merged

pull[bot] merged 53 commits into
Cache-Cloud:masterfrom
tensorflow:master

Conversation

@pull

@pull pull Bot commented Sep 26, 2026 •

Copy link
Copy Markdown

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 : )

Ashutosh0x and others added 30 commits July 29, 2026 23:02
…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.
…mSessionOptions

Adds NumThreadsNegative in process_util_test.cc to verify that
NumInterOpThreadsFromSessionOptions returns a positive thread count
when inter_op_parallelism_threads is negative.

Follow-up to #124954 and commit 1d31888.
Add checks to ensure height and width sequence tensor values are non-negative.
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 &quot;identity&quot; 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 &quot;identity&quot; 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 />

[![Dependabot compatibility score](https://dependabot-badges.githubapp.com/badges/compatibility_score?dependency-name=astral-sh/setup-uv&package-manager=github_actions&previous-version=10.0.1&new-version=10.1.0)](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
PiperOrigin-RevId: 987427673
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
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
EusebioDM and others added 23 commits September 25, 2026 09:38
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
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
…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
Also includes a quant_spec metadata in tfl.fully_connected to allow encoding custom specs.

PiperOrigin-RevId: 987624321
…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
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
@pull pull Bot locked and limited conversation to collaborators Sep 26, 2026
@pull pull Bot added the ⤵️ pull label Sep 26, 2026
@pull
pull Bot merged commit b5d8801 into Cache-Cloud:master Sep 26, 2026
2 checks passed
Sign up for free to subscribe to this conversation on GitHub. Already have an account? Sign in.

Projects

None yet

Development

Successfully merging this pull request may close these issues.