Repository navigation
[pull] master from tensorflow:master - #8837
Merged
Merged
Conversation
TopK used a 32-bit loop index. Adding the thread stride near INT32_MAX could overflow the index and make the kernel read outside the tensor. I changed the loop index to int64. It is converted back to int when TopK stores the result because the output indices are int32.
Co-authored-by: gemini-code-assist[bot] <176961590+gemini-code-assist[bot]@users.noreply.github.com>
kaiser_window called bessel_i0e with abs(beta). Since the second derivative of abs vanishes at the origin, the resulting cusp made the second derivative of the window with respect to beta evaluate to zero at beta == 0, even though the window is I0(beta * s) / I0(beta), an even analytic function of beta. bessel_i0e is an even function, so beta is now passed to it with its original sign. The exponential scaling factor keeps using abs(beta), so negative beta values still produce the same result as positive ones and the computation stays numerically stable for large beta. Adds a unit test for the second derivative at beta == 0. Verified that the forward values are bit-for-bit unchanged for beta >= 0 in float16/float32/ float64, and that negative beta now matches positive beta exactly. Fixes #127245
A tf.TensorArray created with dynamic_size=True lowers to TensorListSetItem
with resize_if_index_out_of_bounds set, and the non-XLA kernel grows the
list when the index is past the end. The tf2xla lowering never read that
attribute, so under jit_compile=True the write was silently discarded and
the program returned a truncated result with no error:
ta = tf.TensorArray(tf.bool, size=0, dynamic_size=True)
# writes at 0 and 1 inside a while_loop, then stack()
eager -> [True, True]
tf.function -> [True, True]
tf.function(jit_compile=True) -> []
The loss is not specific to size=0. XLA freezes the buffer at the size
TensorListReserve allocated and drops every write beyond it, so size=1 with
two writes returns [True] just as quietly. ExecuteTensorListSetItem lowers
to a DynamicUpdateSlice, which clamps its start indices rather than
failing, so the out-of-bounds element lands back inside the buffer.
XLA cannot support the growth itself, since that needs a reallocation and
the buffer shape is static, and it cannot bounds check at runtime either,
since a compiled program cannot raise a data-dependent error. Reject the
op at compile time instead, so the failure is a clear message rather than
a wrong answer.
This also makes the two lowerings of the same user-level API agree. The
older TensorArray path already rejects dynamic_size unconditionally in
TensorArrayOp in tensor_array_ops.cc; only the TensorList path that TF2
actually emits was missing the check.
Fixes #127528
The scaling factor was computed as `beta_abs * shape - beta_abs`. Near the centre of the window `shape` is close to one, so the two products nearly cancel and the difference loses precision. For large `beta` the result can round to exactly zero, making the factor evaluate to 1.0 instead of the correct value. Factor `beta_abs` out so that the subtraction is performed first: `beta_abs * (shape - 1)`. The two forms are algebraically identical and differ only in floating-point rounding. Measured against a higher precision reference, the new form is the more accurate one, by up to ~400x in float64 for large `beta`. Forward values therefore change slightly, by at most ~1.6e-2 relative in float16 at `beta = 60`, ~2.0e-6 in float32 and ~3.5e-15 in float64, where the new values are closer to the exact result. The second derivative at `beta == 0` and the treatment of negative `beta` are unaffected.
The comment credited the loss entirely to DynamicUpdateSlice clamping, but that only covers the case where the element fits and the index is past the end. When the element does not fit at all, which is what a list reserved with size 0 hits, ExecuteTensorListSetItem takes an earlier return that logs a warning and returns the list unchanged with an OK status. Describe both, since the reported repro goes through the first one. No functional change.
Uses tensor_list rather than l in the two tests added here, per review feedback. The rest of this file uses l under its g-bad-name pragma, so this leaves the new tests inconsistent with their neighbours; flagging that for the reviewer to settle either way. Also corrects the same inaccuracy fixed in the kernel comment: the write is dropped either by the early return when the element does not fit, or by DynamicUpdateSlice clamping an out-of-range index, not by clamping alone. No functional change; both tests were re-checked against an unpatched 2.21.0 wheel and still fail there as intended.
The classic tf2xla lowering is not the only path that compiles these ops. When the MLIR bridge runs, CreateTensorListOpsDecompositionPass is added in compile_mlir_util.cc under lower_to_xla_hlo, and HandleTensorListSetItemOp decomposed the op without reading resize_if_index_out_of_bounds, so the same write was dropped just as quietly there. Reject it in the pass too, with the same message the kernel uses, and add a lit case covering it. The buffer the pass builds has a static shape for the same reason the kernel's does, and cutil::SetElement lowers to a DynamicUpdateSlice that clamps an out-of-range index rather than failing. The Python tests now accept InvalidArgumentError as well as UnimplementedError, since emitOpError surfaces as the former, and pin the assertion on the message instead of the type so either path satisfies it.
The new tests wrapped their bodies in def_function.function with jit_compile=True and ran them under self.session(). In that harness the ops are instantiated as ordinary CPU kernels and XlaCompiler never runs, so neither rejection fired and both tests failed with the error not raised. Switch them to the idiom the rest of this file uses, building the ops under self.session() with self.test_scope() and evaluating the result. That is how testInvalidSplitLength asserts UnimplementedError out of the same lowering, and how testGetItemFromEmptyList reaches these kernels. No test in tensorflow/compiler/tests asserts a compile error through jit_compile, which is the hint that the previous shape was wrong. Drops the end to end TensorArray test. There is no established way here to make a dynamic_size TensorArray in a while_loop reach the XLA kernel under this harness, and a test that cannot reach the code it covers is worse than no test. The op level test covers the kernel rejection and the new lit case covers the MLIR one. Removes the two imports and the two BUILD deps that test had needed.
Per review: making this a hard compile error rejects every graph that sets resize_if_index_out_of_bounds, which is every dynamic_size TensorArray, including the many that never actually write past the end and compile correctly today. That is too broad to land. Warn instead and let compilation proceed: - the tf2xla kernel logs the message rather than failing construction - the MLIR pass uses emitWarning rather than emitOpError, so decomposition continues - the lit case expects a warning - the Python test evaluates the op and checks the result instead of asserting a raise, with a counterpart pinning that the attribute alone changes nothing but the logging Note this leaves the truncation itself in place, so the wrong result reported in the issue still occurs; the change only adds a log trail. Adds the logging.h include the kernel now needs.
…derivative PiperOrigin-RevId: 983795354
PiperOrigin-RevId: 983800013
Imported from GitHub PR openxla/xla#48926 Bumps [github/codeql-action/upload-sarif](https://github.com/github/codeql-action) from 4.37.9 to 4.38.0. <details> <summary>Release notes</summary> <p><em>Sourced from <a href="https://github.com/github/codeql-action/releases">github/codeql-action/upload-sarif's releases</a>.</em></p> <blockquote> <h2>v4.38.0</h2> <ul> <li>On GitHub-hosted runners, the CodeQL Action now deletes unused CodeQL bundles from the toolcache before downloading a different bundle, which frees up disk space for the analysis. We expect to roll this change out to everyone in September. <a href="https://redirect.github.com/github/codeql-action/pull/4124">#4124</a></li> <li>The CodeQL Action now supports CodeQL releases that are compatible with Linux Arm64 and downloads the native <code>linux-arm64</code> CodeQL bundle when available. <a href="https://redirect.github.com/github/codeql-action/pull/4072">#4072</a></li> <li>Update default CodeQL bundle version to <a href="https://github.com/github/codeql-action/releases/tag/codeql-bundle-v2.27.0">2.27.0</a>. <a href="https://redirect.github.com/github/codeql-action/pull/4129">#4129</a></li> </ul> </blockquote> </details> <details> <summary>Changelog</summary> <p><em>Sourced from <a href="https://github.com/github/codeql-action/blob/main/CHANGELOG.md">github/codeql-action/upload-sarif's changelog</a>.</em></p> <blockquote> <h1>CodeQL Action Changelog</h1> <p>See the <a href="https://github.com/github/codeql-action/releases">releases page</a> for the relevant changes to the CodeQL CLI and language packs.</p> <h2>[UNRELEASED]</h2> <p>No user facing changes.</p> <h2>4.38.0 - 09 Sept 2026</h2> <ul> <li>On GitHub-hosted runners, the CodeQL Action now deletes unused CodeQL bundles from the toolcache before downloading a different bundle, which frees up disk space for the analysis. We expect to roll this change out to everyone in September. <a href="https://redirect.github.com/github/codeql-action/pull/4124">#4124</a></li> <li>The CodeQL Action now supports CodeQL releases that are compatible with Linux Arm64 and downloads the native <code>linux-arm64</code> CodeQL bundle when available. <a href="https://redirect.github.com/github/codeql-action/pull/4072">#4072</a></li> <li>Update default CodeQL bundle version to <a href="https://github.com/github/codeql-action/releases/tag/codeql-bundle-v2.27.0">2.27.0</a>. <a href="https://redirect.github.com/github/codeql-action/pull/4129">#4129</a></li> </ul> <h2>4.37.9 - 26 Aug 2026</h2> <ul> <li>Update default CodeQL bundle version to <a href="https://github.com/github/codeql-action/releases/tag/codeql-bundle-v2.26.4">2.26.4</a>. <a href="https://redirect.github.com/github/codeql-action/pull/4106">#4106</a></li> </ul> <h2>4.37.8 - 21 Aug 2026</h2> <p>No user facing changes.</p> <h2>4.37.7 - 13 Aug 2026</h2> <ul> <li>Update default CodeQL bundle version to <a href="https://github.com/github/codeql-action/releases/tag/codeql-bundle-v2.26.3">2.26.3</a>. <a href="https://redirect.github.com/github/codeql-action/pull/4085">#4085</a></li> </ul> <h2>4.37.6 - 04 Aug 2026</h2> <ul> <li>Changed the default filepath for the new remote file address format that was introduced in CodeQL Action 4.37.0 / 3.37.0 to <code>.github/codeql-config.yml</code> to align it with the suggested path that is used elsewhere. <a href="https://redirect.github.com/github/codeql-action/pull/4070">#4070</a></li> </ul> <h2>4.37.5 - 03 Aug 2026</h2> <ul> <li>Fixed a bug where a network error while streaming the download of the CodeQL bundle could terminate the <code>init</code> Action instead of falling back to downloading the bundle before extracting it. <a href="https://redirect.github.com/github/codeql-action/pull/4061">#4061</a></li> </ul> <h2>4.37.4 - 29 Jul 2026</h2> <ul> <li>This version of the CodeQL Action adds support for the <code>tools</code> input for the <code>codeql-action/init</code> step to be specified using a <code>github-codeql-tools</code> <a href="https://docs.github.com/en/organizations/managing-organization-settings/managing-custom-properties-for-repositories-in-your-organization">repository property</a>. This feature will gradually be rolled out following the release of this version. Once rolled out, this allows for the CodeQL CLI version that is used in GitHub-managed workflows, such as Default Setup, to be set to a custom value. For example, customers who run into issues with rate limits when a new CodeQL CLI version is released can set the value to <code>toolcache</code> to always use the CodeQL CLI version that is available in the runner toolcache. For Advanced Setup workflows, the value provided for <code>tools</code> in the workflow definition always takes precedence unless the value of the repository property starts with <code>!</code>. <a href="https://redirect.github.com/github/codeql-action/pull/4037">#4037</a></li> <li>Update default CodeQL bundle version to <a href="https://github.com/github/codeql-action/releases/tag/codeql-bundle-v2.26.2">2.26.2</a>. <a href="https://redirect.github.com/github/codeql-action/pull/4051">#4051</a></li> </ul> <h2>4.37.3 - 22 Jul 2026</h2> <p>No user facing changes.</p> <h2>4.37.2 - 21 Jul 2026</h2> <ul> <li>The new address format for the <code>config-file</code> input that was introduced in CodeQL Action 4.37.0 is now enabled by default. In addition to the format described there, the <code>remote=</code> prefix can now be used to explicitly indicate that the input refers to a remote file. All previous input formats continue to be accepted as well. <a href="https://redirect.github.com/github/codeql-action/pull/4023">#4023</a></li> <li>The CodeQL Action can now make use of <a href="https://docs.github.com/en/code-security/how-tos/secure-at-scale/configure-organization-security/manage-usage-and-access/giving-org-access-private-registries">configured private registries</a> in Default Setup to retrieve CodeQL configuration files from remote repositories that require authentication. This will allow customers to store their CodeQL configuration in a single repository that can then be referenced by Default Setup workflows in other repositories. We expect to roll this and other, related changes out to everyone in July. <a href="https://redirect.github.com/github/codeql-action/pull/4007">#4007</a></li> </ul> <h2>4.37.1 - 16 Jul 2026</h2> <!-- raw HTML omitted --> </blockquote> <p>... (truncated)</p> </details> <details> <summary>Commits</summary> <ul> <li><a href="https://github.com/github/codeql-action/commit/b96794f015dfd88f77b49b1c93e0fa7110f94c63"><code>b96794f</code></a> Merge pull request <a href="https://redirect.github.com/github/codeql-action/issues/4131">#4131</a> from github/update-v4.38.0-7e08580a9</li> <li><a href="https://github.com/github/codeql-action/commit/02d5093871674ea20274117103ce3038c73c77ef"><code>02d5093</code></a> Update changelog for v4.38.0</li> <li><a href="https://github.com/github/codeql-action/commit/7e08580a93dc4e4b9dda167e364577035cf504c6"><code>7e08580</code></a> Merge pull request <a href="https://redirect.github.com/github/codeql-action/issues/4130">#4130</a> from github/henrymercer/workflow-runner-sizing</li> <li><a href="https://github.com/github/codeql-action/commit/bfcc52b4f5d98468a5993daa0bf0e4fb3f3ed698"><code>bfcc52b</code></a> Run slow macOS checks on larger runners</li> <li><a href="https://github.com/github/codeql-action/commit/8c251e757c0260283fc50214a06ac768b61d3af4"><code>8c251e7</code></a> Merge pull request <a href="https://redirect.github.com/github/codeql-action/issues/4129">#4129</a> from github/update-bundle/codeql-bundle-v2.27.0</li> <li><a href="https://github.com/github/codeql-action/commit/0b7ca400df35985869d4b9146a067865d4115da1"><code>0b7ca40</code></a> Add changelog note</li> <li><a href="https://github.com/github/codeql-action/commit/40484b339517c6bcf00f81eebc95ca041ddca505"><code>40484b3</code></a> Update default bundle to codeql-bundle-v2.27.0</li> <li><a href="https://github.com/github/codeql-action/commit/977e6ceaea7361825998245d787fa3b4d6b9e5df"><code>977e6ce</code></a> Merge pull request <a href="https://redirect.github.com/github/codeql-action/issues/4124">#4124</a> from github/henrymercer/toolcache-bundle-cleanup</li> <li><a href="https://github.com/github/codeql-action/commit/40a6b3824794ae1156e1a5320d32e364bf1dcebc"><code>40a6b38</code></a> Address toolcache cleanup review feedback</li> <li><a href="https://github.com/github/codeql-action/commit/deece8f852f048bc3f52fd42c9cc7a99b1ebb252"><code>deece8f</code></a> Apply suggestion from <a href="https://github.com/henrymercer"><code>@henrymercer</code></a></li> <li>Additional commits viewable in <a href="https://github.com/github/codeql-action/compare/cdf488f595d80d6e07e03d4674febd5ab45fa938...b96794f015dfd88f77b49b1c93e0fa7110f94c63">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: -- 64555e142a81a3f779a0078e93b9e597d04712dd by dependabot[bot] <49699333+dependabot[bot]@users.noreply.github.com>: Bump github/codeql-action/upload-sarif from 4.37.9 to 4.38.0 Bumps [github/codeql-action/upload-sarif](https://github.com/github/codeql-action) from 4.37.9 to 4.38.0. - [Release notes](https://github.com/github/codeql-action/releases) - [Changelog](https://github.com/github/codeql-action/blob/main/CHANGELOG.md) - [Commits](github/codeql-action@cdf488f...b96794f) --- updated-dependencies: - dependency-name: github/codeql-action/upload-sarif dependency-version: 4.38.0 dependency-type: direct:production update-type: version-update:semver-minor ... Signed-off-by: dependabot[bot] <support@github.com> Merging this change closes #48926 PiperOrigin-RevId: 983804916
PiperOrigin-RevId: 983805458
With my previous DWYU change I accidentally registered a loop in the macro expansion which makes DWYU segfault. This is the fix for that. Also improves testing. PiperOrigin-RevId: 983805520
Imported from GitHub PR openxla/xla#48981 📝 Summary of Changes Skip NumericTestsForBlas.Infinity/dot_tf32_tf32_f32_x3 on MI350 (gfx950). 🎯 Justification hipBLASLt emulates TF32 on gfx950 and returns NaN for inf * 1. TF32x3 relies on that high×high product after zeroing the low order terms, so the test fails on MI350. It still passes on gfx942 and with Triton, so the skip is scoped to gfx950 only. 🚀 Kind of Contribution Please remove what does not apply: 🧪 Tests 🧪 Unit Tests: GTEST_SKIP() on ROCm gfx950: xla/backends/gpu/codegen/triton/dot_algorithms_test.cc — NumericTestsForBlas/NumericTestsForBlas.Infinity/dot_tf32_tf32_f32_x3 Copybara import of the project: -- 6e41451edc164f240c9bed076324a5adfed7b6cb by Zoran Jovanovic <zjovanov@amd.com>: [ROCm] Skip NumericTestsForBlas.Infinity TF32x3 on MI350 Merging this change closes #48981 PiperOrigin-RevId: 983805760
Imported from GitHub PR openxla/xla#48110 ## 📝 Summary of Changes Replace the compile-time `TF_ROCM_VERSION` gate on the C64/C128 rocBLAS redirect from #44047 with a runtime capability check. Complex matmuls take the hipBLASLt plan when the descriptor builds and the heuristic offers an algorithm, and fall back to rocBLAS (`rocblas_cgemm`/`zgemm`) otherwise. ## 🎯 Justification hipBLASLt can support complex C64/C128 GEMM kernels in ROCm 10, but a build-time macro describes the ROCm that XLA was compiled against, not the hipBLASLt loaded at runtime. so making the check to runtime so legacy rocms will still use rocblas and newer one switch to hipblasLT. ## 🚀 Kind of Contribution ♻️ Cleanup Copybara import of the project: -- 8e63c8691ee87cd538cf16ec811faeb5bbbb56f7 by magaonka <magaonka@amd.com>: [ROCm] Pick the rocBLAS complex GEMM fallback at runtime hipBLASLt gained complex GEMM kernels in ROCm 10, but whether they are usable depends on the hipBLASLt loaded at runtime and on the device, not on the ROCm version XLA was built against. A compile-time TF_ROCM_VERSION gate strands mismatched build/runtime combinations on the wrong path in both directions. Try the hipBLASLt plan for C64/C128 and keep it only if the descriptor builds and the heuristic offers an algorithm; otherwise redirect to the openxla/xla#44047 rocBLAS path. Both conditions are needed: older hipBLASLt rejects the complex matmul descriptor outright, while a newer one could still be missing Tensile kernels for a given architecture. Merging this change closes #48110 PiperOrigin-RevId: 983807962
…enchmark_test on ROCm Imported from GitHub PR openxla/xla#48711 📝 Summary of Changes - Enable `local_client_execute_test` on ROCm by dropping the stale `cuda-only` tag (keep `optonly`). - Enable `gpu_bandwidth_benchmark_test` on ROCm by linking `rocm_platform` and switching to `xla_test`. - Keep NVIDIA on H100 (`requires-gpu-sm90-only`) and add `amdgpu_any` for ROCm. 🎯 Justification Add more unit test coverage on ROCm platform. 🚀 Kind of Contribution 🧪 Tests Copybara import of the project: -- 3686a353cc8bc068e29da882093b7e5bb435ae04 by linchen1 <lin.chen1@amd.com>: Enable local_client_execute_test and gpu_bandwidth_benchmark_test on ROCm. Signed-off-by: linchen1 <lin.chen1@amd.com> Merging this change closes #48711 PiperOrigin-RevId: 983808132
…red address indexes Imported from GitHub PR openxla/xla#48789 ## 📝 Summary of Changes Reduce the host overhead of device_address_vmm_allocator with below changes: 1. Replace the VMM allocator's address hash maps with ordered `absl::btree_map` indexes for targeted overlap checks. 2. Store deferred deallocations in an intrusive queue so reused allocations can cancel their pending frees directly, while preserving mapping reuse and deferred teardown behavior. ## 🎯 Justification Avoid scanning all tracked ranges during overlap checks and searching the pending queue when reusing a known allocation. This reduces CPU bookkeeping for workloads with many live or deferred VMM allocations. ## 🚀 Kind of Contribution ⚡️ Performance Improvement, 🧪 Tests ## 🧪 Unit Tests: Added coverage in `//xla/stream_executor:device_address_vmm_allocator_test`: - `OverlapLookupPreservesRangeBoundariesForActiveAndStaleAddresses` - `ManyReservationAliasesReuseAndRemoveOrderedRecords` - `StaleAliasReuseValidatesSourceAndFullReservationRange` - `MixedPendingOperationsCanBeCancelledAndRequeuedInAnyOrder` Copybara import of the project: -- 44cc7cb82c2c704a78e91cf641530fa7b3391ab2 by Shawn Wang <shawnw@nvidia.com>: Optimize VMM allocator overlap lookups with ordered address indexes -- cf69d2304c5bb34736c6c51501d110f931fbc557 by Shawn Wang <shawnw@nvidia.com>: Reuse VMM aliases before overlap checks and unlink pending operations Reactivate an exact stale reservation alias directly from its validated source record before searching either destination address index. Preserve the existing validation and waiting path for new or changed mappings. Embed intrusive pending-operation nodes in allocation and alias records. Reuse now cancels the known operation in constant time without scanning or shifting queued entries, or allocating separate queue nodes. Preserve FIFO stream ordering, batch accounting, and deferred teardown. Add tests for full-range/source validation, repeated cancellation and requeueing of mixed operations, and draining uncancelled aliases. Validation: 11 CPU allocator tests and 34 CUDA allocator tests passed, including the two-device cases with CUDA_VISIBLE_DEVICES=0,1. -- 740911322f61c8ccebade5106ff542c299620846 by Shawn Wang <shawnw@nvidia.com>: Simplify VMM allocator reuse and deferred completion Build deferred-operation keys from queue nodes and centralize unlinking and metadata snapshots in completion. Remove the redundant address-index lookup when reactivating an allocation and flatten the stale-alias branch made redundant by early reuse. Preserve the early alias reuse path, constant-time queue cancellation, and existing validation and synchronization behavior. Validation: 11 CPU and 34 CUDA allocator tests passed with no skips in the maxtext-2026-09-12 container, using at most two B200 GPUs. Merging this change closes #48789 PiperOrigin-RevId: 983809891
A packing spec's `AddAddressArgument(i)` indices refer to positions in the device buffer list that the runtime builds from `emitters::KernelArguments`. Nothing checks that those indices are in range until `BuildArguments` runs, at which point the kernel is already being launched, so an off-by-one between a kernel's signature and the emitter that set up its arguments only shows up as a runtime error on the machine executing the program. `Validate(num_available_args)` performs that check up front, so emitters can fail at compile time with a message that points at the emitter rather than at the launch. `size()`, `MaxArgumentIndex()` and `KernelArgPackingSpec::relocation_argument_index()` expose the information it needs and are useful for diagnostics on their own. `Identity(num_args)` covers the common "pass every buffer pointer straight through, in order" case. `ptx_custom_kernel.cc` had a private copy of it; it now uses the shared one. PiperOrigin-RevId: 983818570
…nsorFlow - Bump rules_cc from 0.2.20 to 0.2.25 in MODULE.bazel and workspace3.bzl for both XLA and TensorFlow. - Bump bazel_features from 1.25.0 to 1.50.0 in workspace3.bzl for both XLA and TensorFlow to provide bazel_features.cc.cc_common_is_in_rules_cc used by rules_cc 0.2.25 in WORKSPACE mode. PiperOrigin-RevId: 983818634
…ot_algorithms_test`. `dot_algorithms_test.cc` directly includes `absl/strings/string_view.h`, which DWYU (`bant`) requires as an explicit dependency now that `xla_test` targets are checked. PiperOrigin-RevId: 983835500
…er and report argument positions. Two changes that let out-of-tree emitters (in particular XLA:GPU native custom-call handlers) build `KernelArguments` correctly instead of assembling `KernelArgument`s by hand: 1. A `Create` overload that takes a `SliceProvider` callback instead of a `BufferAssignment`. `ThunkEmitter` resolves buffers through `GetAllocationSlice`, which first consults per-instruction allocation overrides and only then falls back to the buffer assignment. Emitters with such overrides cannot use the `BufferAssignment` overloads without silently dropping the override and launching the kernel on the wrong buffers. 2. `OperandIndex` and `ResultIndex`, which report where an operand or a result subshape ended up in `args()`. Those positions are the index space of `stream_executor::KernelArgsPackingSpec` relocations, so callers previously had to hardcode them and keep them in sync with the argument layout by hand. This is particularly error-prone for the interleaved overload. Both return `FailedPrecondition` when the arguments were assembled by hand rather than by a `Create` factory. All overloads now share a single private `CreateInternal` implementation, including the one that interleaves results with operands. That overload used to duplicate the construction logic, and had already drifted from it: it ignored the `xla.no_invariant_operands` frontend attribute, so an operand the producer marked as non-invariant was passed to the kernel as invariant anyway. It now also rejects interleaving indices that are out of order or that do not cover every result, rather than reporting a confusing count mismatch afterwards. PiperOrigin-RevId: 983836076
…lternative_layout_test on ROCm Imported from GitHub PR openxla/xla#48531 📝 Summary of Changes Drop cuda-only from both convolution_1d_test and convolution_1d_gpu_alternative_layout_test to enable them on ROCm CI. 🎯 Justification Add more unit test coverage on ROCm platform. 🚀 Kind of Contribution 🧪 Tests Copybara import of the project: -- c703f27f257495c05af87b7edcb971abadc9626a by Lin Chen1 <lin.chen1@amd.com>: Enable convolution_1d_test and convolution_1d_gpu_alternative_layout_test on ROCm. Signed-off-by: Lin Chen1 <lin.chen1@amd.com> Merging this change closes #48531 PiperOrigin-RevId: 983836641
PiperOrigin-RevId: 983845225
Imported from GitHub PR openxla/xla#48416 This PR temporarily updates the PJRT GPU plugin name to SYCL to pass unit tests, we are reverting back to naming the PLUGIN NAME as SYCL to be consistent with the stream executor name of SYCL. Unit Test Fixes: //xla/pjrt/c/c:pjrt_c_api_gpu_test_intelgpu_any Copybara import of the project: -- a07c441910f7fdb60c7b4b6b994b8b9162147a66 by Ashiq Imran <ashiq.imran@intel.com>: fix for PJRT C API -- d51264f8e3d72ef99d48ae8017a4e806f3a9c1c5 by Ashiq Imran <ashiq.imran@intel.com>: minor changes Merging this change closes #48416 PiperOrigin-RevId: 983845645
…back: 1. Support CUDA 12 and CUDA 13 pip extras in setup.py (tensorflow[cuda12], tensorflow[cuda13]). 2. Dynamically resolve and fallback to alternative CUDA component DSO versions at runtime. 3. Configure Kokoro release jobs for CUDA 13 wheels. PiperOrigin-RevId: 983851549
Splits CompositeRewriter::RewriteComputation into modular precondition helper functions returning xla::Decision instead of boolean flags, and extracts instruction rewrite logic into RewriteScaledDot. PiperOrigin-RevId: 983862544
Imported from GitHub PR openxla/xla#48533 ## Summary of Changes Reject BFC allocation requests that cannot be rounded up to the minimum allocation size without overflowing size_t. Add a focused allocator regression test for the overflow boundary. ## Justification A request near SIZE_MAX currently wraps to zero during rounding. The allocator can then return an undersized chunk and later fail while validating corrupted chunk metadata. This was observed through TensorFlow in #108904, but the affected allocator is maintained in OpenXLA. ## Kind of Contribution Bug fix and tests. ## Benchmark Not applicable; this only adds a cold-path bounds check before allocation. ## Unit Tests The new BFCAllocatorTest.RejectsAllocationWhoseRoundedSizeWouldOverflow test uses the fake suballocator and does not reserve a large buffer. Local validation: git diff --check. ## Execution Tests Not applicable; the regression is covered at the allocator boundary. Copybara import of the project: -- 59caf3bfacc5771f65333d5e0a0701519bc54e4f by Sai Ramineni <csramineni@gmail.com>: Prevent overflow when rounding BFC allocation sizes Merging this change closes #48533 PiperOrigin-RevId: 983868621
Imported from GitHub PR openxla/xla#47777 📝 **Summary of Changes** Modified the `RaggedAllToAll` GPU collective in `xla/backends/gpu/runtime/ragged_all_to_all_thunk.cc` to explicitly skip `LaunchSend`, `LaunchRecv`, and `LaunchPut` calls for zero-size peer slices in the NCCL fallback path. In the Put+Signal path, it now tracks the actual number of puts issued per peer and passes this correct `op_cnt` to `WaitSignal`. 🎯 **Justification** This change resolves issue #46982 by eliminating an O(n_devices) cost scaling bottleneck during ragged-all-to-all operations with sparse neighbor exchange patterns. Previously, NCCL enqueued work for zero-count operations, resulting in ~12 µs overhead for every zero-size slot. By skipping these slots, this change significantly improves performance for sparse communication on a large number of GPUs (e.g., eliminating the +625 µs scaling overhead from 16 to 64 devices on a fixed degree-10 neighbor graph). 🚀 **Kind of Contribution** 🐛 Bug Fix, ⚡️ Performance Improvement 📊 **Benchmark (for Performance Improvements)** Eliminates the ~12 µs per zero-size slot NCCL overhead reported in #46982 (cost is now constant O(actual_neighbors) rather than linear O(n_devices)). 🧪 **Unit Tests:** Existing tests in `ragged_all_to_all_thunk_test.cc` continue to validate thunk initialization and command buffer logic. 🧪 **Execution Tests:** Added `RaggedAllToAll_2GPUs_SparseNeighbors` end-to-end test to `xla/backends/gpu/tests/ragged_all_to_all_e2e_test.cc`. This 2-GPU test verifies the collective's correctness under a sparse pattern where one replica sends zero rows in one direction, ensuring that skipping zero-size slots works correctly without deadlocks and produces the expected data. Copybara import of the project: -- f90c7f3f9567c81587085eaac28ab5f34182a0e3 by ankit <ankit99aug@gmail.com>: Fix O(n_devices) overhead in GPU ragged-all-to-all Fixes #46982 by skipping zero-size peer slices in the NCCL fallback path. -- bc5ec206d7881d118da64513c4cfd03d23705543 by Ankit Vishwakarma <ankit99aug@gmail.com>: Fix zero-size ragged all-to-all operations Merging this change closes #47777 PiperOrigin-RevId: 983874052
- Added positive and negative type inference tests to ops_chlo.mlir and a test to chlo_legalize_to_stablehlo.mlir. - Made some changes in ChloLegalizeToStablehlo.cpp so that the bounds are preserved for dynamic tensors. PiperOrigin-RevId: 983880829
PiperOrigin-RevId: 983896161
…g runtimes. Eliminates redundant in-memory copies of GraphDef and FunctionLibraryDefinition during model initialization and signature compilation in TensorFlow serving runtimes. PiperOrigin-RevId: 983901136
…te-stride span Imported from GitHub PR openxla/xla#48513 #ai-generated I found this while reviewing the recently landed #46604 byte-stride overflow fix and verified the remaining accumulation wrap against the current XLA tree with an in-tree Bazel regression test. `FromZerothElementPointer` already rejects a `stride * (dim - 1)` product that overflows `int64`. Adding those products into a `uint64_t` without an overflow check still wraps: ```text shape = [2, 2, 2] strides = [INT64_MAX, INT64_MAX, 2] dtype = u8 accepted region size = 1 ``` A second case, `strides = [INT64_MAX, INT64_MAX, 1]`, accumulates to `UINT64_MAX` and only overflows on the final `+ dtype-size` add. This change checks `__builtin_add_overflow` on the running sum and on the final add. Tests: `ByteStrideAccumOverflowUint64`, `ByteStrideFinalSizeOverflowUint64`, plus compact `u8` `[2,2,2]` success cases. Compact-layout `byte_size * shape.num_elements()` is not changed (#48286). ```text bazel test //xla/python/ifrt_proxy/common:array_util_test ``` **PASSED**, 42/42, including `ByteStrideAccumOverflowUint64` and `ByteStrideFinalSizeOverflowUint64`. Copybara import of the project: -- 4116a8f9c0eda07ba8b6d1ffe36ba22bc9541045 by roopd3v <roopd3v@users.noreply.github.com>: Fix uint64 accumulation overflow in IFRT ArrayMemRegion byte-stride span. The per-term stride*(dim-1) multiply is already checked. Adding those products into a uint64_t without an overflow check still wraps, so a client-supplied stride vector such as [INT64_MAX, INT64_MAX, 2] with shape [2, 2, 2] is accepted as a 1-byte region. Check the running sum and the final + dtype-size add. Tests cover both the accumulation wrap and the final add (INT64_MAX + INT64_MAX + 1 == UINT64_MAX). Merging this change closes #48513 PiperOrigin-RevId: 983915659
…ously constructed `BufferAllocationProto::Assigned` protobuf objects in hash-map order and then sorted the repeated protobuf field by moving/swapping protobuf messages. In addition, `BufferAssignment::ToProto` constructed temporary stack `LogicalBufferProto` and `LogicalBufferProto::Location` messages for every logical buffer and alias before swapping them into the parent message. PiperOrigin-RevId: 983932289
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.
19 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.
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 : )