Repository navigation
[pull] master from tensorflow:master - #8852
Merged
Merged
Conversation
TensorListScatter resized with signed `highest_index + 1`, which overflows when the index is INT32_MAX. Apply the same guard already used by TensorListScatterIntoExistingList and size the backing vector with size_t.
The registered gradient for ImageProjectiveTransformV2/V3 inverted the transform matrix and re-ran the forward resampling op on the incoming gradient. That is only a correct adjoint when the forward coordinate map is a bijection on the pixel grid -- never true once fill_mode clamps or folds an out-of-range coordinate, or the transform downscales. In both cases it silently duplicates or drops gradient mass at the affected pixels. Replace it with a real scatter-based adjoint: for each output pixel, scatter-add its incoming gradient onto the same input pixel(s) the forward kernel actually read, with the same interpolation weights. The coordinate mapping (_map_coordinate) mirrors MapCoordinate in tensorflow/core/kernels/image/image_ops.h exactly, including its use of C++ truncation-toward-zero and round-half-away-from-zero, which TensorFlow's own division/round otherwise disagree with. Fixes #127241, where tf.keras.ops.image.affine_transform with interpolation="nearest" returned gradient 30.0 for a case where the correct value is 17.0. Verified against a real build (TF 2.21.0, matching the report): - The issue's exact reproducer now returns 17.0. - Adjoint identity <forward(x), g> == <x, grad_fn(g)> holds (rel. error <= 1.9e-6) across all 4 fill_modes x 2 interpolations x 6 transform types x batch/broadcast x float32/float64, 384 configs. - _map_coordinate matches the C++ kernel's actual pixel selection on 4420 probed pixels across fill_modes and transforms. - gradient_checker_v2 agrees with finite differences to <8e-5. - pylint 10.00/10 under the repo's pylintrc (Python 3.9, its CI config). Adds ImageProjectiveTransformOpTestBase covering all of the above as regression tests; every one of them fails against the old gradient and passes against the new one.
- Guard _map_coordinate's WRAP/REFLECT division against length == 1 (a 1-pixel-wide/tall axis), which otherwise divides by zero. The masked-out result was already correct either way, but the intermediate inf/nan getting cast to int32 is platform-dependent undefined behavior, worth avoiding regardless of whether its value is later discarded. - Build batch_idx via array_ops.broadcast_to instead of a cast-to-float + zeros_like(x) + add round trip, and keep it int32 throughout, dropping the now-redundant cast where it's used. - Use array_ops.scatter_nd directly instead of zeros + tensor_scatter_add; scatter_nd sums duplicate indices the same way, verified directly against tf.scatter_nd, so this is behavior-preserving for the overlapping-tap case this PR exists to fix. Drops the now-unused gen_array_ops import. - Add testOnePixelWideInputIsFinite covering the length == 1 case the first fix addresses. Re-verified after applying: the issue's reproducer, the 384-config adjoint-identity sweep, the 4420-pixel MapCoordinate parity probe, and gradient_checker_v2 all still pass with identical results; pylint still 10.00/10; full ImageProjectiveTransformOpTest suite (29 tests, 12 new) passes.
Both targets use strict_deps, so direct deps must match direct imports: - image_ops: add :array_ops_stack, :clip_ops and :math_ops (newly imported by the adjoint implementation), drop :linalg_ops (no longer used now that the gradient no longer inverts the transform matrix). - image_grad_test_base: add :array_ops.
- Add Args/Returns sections to the _trunc_div and _round_half_away_from_zero docstrings. - Fix continuation-line indentation in ImageProjectiveTransformOpTestBase._transform. No functional change.
Per review: the internal XLA GPU stack executes the forward op's coordinate mapping with enough numerical divergence from CPU/eager to trip this suite's tight tolerances (e.g. testAdjointIdentity's rtol=atol=1e-6), unrelated to the adjoint logic itself. Follows the same test_util.for_all_test_methods(test_util.disable_xla, ...) pattern already used above for ResizeNearestNeighborOpTestBase for an analogous reason.
PiperOrigin-RevId: 987096919
…2 MiB Compiling GPU convolutions invokes cuDNN v9 engine heuristics (cudnnBackendGetAttribute -> traceback_api_add), which together with XLA GPU compilation frames requires ~280+ KiB of stack space and overflows the previous 240 KiB stack size (breaking jax_api_pjrt_gpu_test after cl/986842563). PiperOrigin-RevId: 987097843
PiperOrigin-RevId: 987118567
… prevent timeout. PiperOrigin-RevId: 987130327
PiperOrigin-RevId: 987131214
`tf.linalg.band_part` instead of raising `InvalidArgumentError`, so that eager execution matches the documented definition of the op and the XLA lowering. Fixes #110316 PiperOrigin-RevId: 987133747
Reverts 4470f2f PiperOrigin-RevId: 987140670
PiperOrigin-RevId: 987147690
PiperOrigin-RevId: 987183957
PiperOrigin-RevId: 987188596
PiperOrigin-RevId: 987195419
…tive-transform-grad-127241 PiperOrigin-RevId: 9872330
When `block_host_until_done` is false, `~ExecutionWatchdogScope()` releases the `HangWatchdog::Guard` asynchronously via `stream_->DoHostCallback(...)`, so the guard and its `on_timeout` and `pre_abort` callbacks outlive `ExecutionWatchdogScope` and `GpuExecutable::ExecuteThunks()`. - Copy `ExecutionTimeoutHandler` by value in `ExecutionWatchdogScope::Arm()` instead of capturing a raw `const GpuExecutableRunOptions*` pointer that can dangle when callers allocate `GpuExecutableRunOptions` on the stack. - Separate `ThunkExecutor::ProgressTracker` from `ThunkExecutor::ScopedProgressTracker` via `std::shared_ptr` so `pre_abort` in `GpuExecutable::ExecuteThunksImpl()` captures `tracker->tracker()` by value while `ScopedProgressTracker` remains a stack-local move-only RAII guard that uninstalls the thread-local pointer when leaving the dispatch scope. PiperOrigin-RevId: 987251708
…GPU crash. PiperOrigin-RevId: 987255629
PiperOrigin-RevId: 987275326
PiperOrigin-RevId: 987276596
…compiler/jit/kernels`
Remove `"@xla//xla/service/gpu:gpu_executable_run_options"` from `XLA_OPS_DEPS` (`xla_ops` and `xla_ops_no_jit_rewrite_registration`) in `tensorflow/compiler/jit/kernels/BUILD`, as `GpuExecutableRunOptions` is not included or used by `xla_ops.{h,cc}`.
PiperOrigin-RevId: 987278014
…n and GPU `backend_configs_cc` dep from `mlir_hlo_to_hlo` `MlirToHloConversionOptions::propagate_bitcast_layouts_to_backend_config` defaults to `false` and is never set or referenced anywhere in the codebase. It was added as a temporary workaround for the legacy MHLO-based XLA:GPU elemental IR emitters (which have since been removed), and was the sole reason the hardware-independent `//xla/hlo/translate/mhlo_to_hlo:mlir_hlo_to_hlo` library included `xla/service/gpu/backend_configs.proto.h` (`xla::gpu::BitcastBackendConfig`) and depended on `//xla/service/gpu:backend_configs_cc`. Remove the dead option, its unreachable `if` branch in `ExportXlaOperator(mhlo::BitcastOp)`, and the `//xla/service/gpu:backend_configs_cc` dependency. PiperOrigin-RevId: 987284482
…m `codegen/tiling` and `codegen/xtile` Remove unused `#include "xla/service/gpu/backend_configs.proto.h"` includes and `//xla/service/gpu:backend_configs_cc` build dependencies from `//xla/codegen/tiling:symbolic_tile_analysis`, `//xla/codegen/tiling/experimental:tiled_hlo`, `//xla/codegen/xtile:block_level_parameters`, and `//xla/codegen/xtile:tiling_from_block_parameters` (left over after `BlockLevelFusionConfig` and `Tile` were migrated to `xtile_config.proto`). PiperOrigin-RevId: 987284532
Extracts scaled dot tests from fusion_emitter_device_test into scaled_dot_device_test to isolate scaled dot test cases and reduce test compilation overhead. PiperOrigin-RevId: 987286750
PiperOrigin-RevId: 987297161
Imported from GitHub PR openxla/xla#48977 📝 Summary of Changes - Add //xla/stream_executor/rocm:rocm_blas_lt_test as the StreamExecutor twin of cuda_blas_lt_test. - Keep cuda-only on the CUDA test. Use GpuPlatformName(), required F32 and S8S32 GEMM, skip S8S8F32 if hipBLASLt returns no algorithms. 🎯 Justification Add more unit test coverage on ROCm platform. 🚀 Kind of Contribution 🧪 Tests Copybara import of the project: -- c9d526efeb59b64ace5a3d6b3d81dea2a38f2350 by Lin Chen1 <lin.chen1@amd.com>: Add new rocm_blas_lt_test for hipBLASLt device GEMM. Signed-off-by: Lin Chen1 <lin.chen1@amd.com> -- 7611d7aed1b43c18a91d4a45efbee51346b8215d by Lin Chen1 <lin.chen1@amd.com>: fix clang-format error in rocm_blas_lt_test. Signed-off-by: Lin Chen1 <lin.chen1@amd.com> -- 4e60098349f79f2403982075a07f624265573b28 by linchen1 <lin.chen1@amd.com>: Add a trailing newline to rocm_blas_lt_test.cc. Signed-off-by: linchen1 <lin.chen1@amd.com> Merging this change closes #48977 PiperOrigin-RevId: 987322640
HLO dynamic-slice and dynamic-update-slice clamp each dimension's start index i to [0, dim_size[i] - slice_size[i]]. Previously, AnalyzeDynamicSlice and VerifySliceOffset computed the 1D byte offset directly from unclamped start indices, relying only on runtime 1D buffer clamping to [0, buffer_size - slice_size]. When a slice is contiguous because all dimensions major to a partially-sliced dimension have slice_size == 1, an out-of-bounds offset on a minor dimension or size-1 dimension could stay within the total buffer bounds and write or read at the wrong element offset. Clamp constant offset operands per-dimension in AnalyzeDynamicSlice and VerifySliceOffset, and for loop-dependent offsets verify that the runtime 1D buffer clamping matches the per-dimension clamped HLO byte offset on every loop iteration. Also preserve the sliced operand shape before walking through bitcasts in DynamicSliceFusion::ResolveParameter so Parameter.shape matches the slice rank and strides in ComputeSliceOffset. Fixes #49120 PiperOrigin-RevId: 987353923
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.
13 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 : )