Repository navigation
[pull] master from tensorflow:master - #8851
Merged
Merged
Conversation
…multi-element tensors - Add OP_REQUIRES validation in WriteScalarSummaryOp::Compute (summary_kernels.cc) to verify that the value tensor contains exactly one element (NumElements() == 1) before delegating to the summary writer. - Return errors::InvalidArgument for empty or multi-element tensors, ensuring Python cleanly raises tf.errors.InvalidArgumentError instead of terminating with a segmentation fault. - Allow single-element tensors of any rank (e.g. shape (), (1,), or (1, 1)) in accordance with established summary op API contracts. - Add comprehensive graph and eager mode unit tests in summary_ops_test.py covering valid scalar and single-element tensor successes alongside empty and multi-element tensor rejections. Fixes #113060.
….concat Fixes #124017 Problem: When tf.rank() is used inside a tf.function with a dynamically-shaped input (i.e. TensorSpec(shape=None, ...)), the static rank is unknown at trace time, so tf.rank(image) produces a 0-D (scalar) Tensor. Building a permutation for tf.transpose with: perm = tf.concat([[num_dims - 1], tf.range(num_dims - 1)], axis=0) triggers TensorFlow's auto-packing conversion on the Python list [num_dims - 1]. Auto-packing calls gen_array_ops.pack (i.e. tf.stack) on the list elements, which is the Pack kernel. On GPU, the Pack kernel rejects 0-D scalar inputs with: InvalidArgumentError: ConcatOp : Can't concatenate scalars (use tf.stack instead) [Op:Pack] name: packed The same pattern works on CPU because the CPU Pack kernel handles scalars correctly. Root cause: _autopacking_helper (in tensorflow/python/ops/array_ops.py) calls gen_array_ops.pack directly when all elements are already Tensor objects. The GPU implementation of the underlying Pack op does not support 0-D scalar inputs. Fix: Detect when all elements being packed are 0-D scalar tensors and, in that case, use reshape(scalar, [1]) + concat_v2 instead of pack. This produces an identical 1-D output tensor and works correctly on both CPU and GPU. The guard is applied in both the eager fast-path and the graph-building code path of _autopacking_helper. Tests: Added testConcatScalarRankDynamic to tensorflow/python/kernel_tests/math_ops/transpose_op_test.py. The test traces a tf.function with TensorSpec(shape=None, ...) to force dynamic rank, then constructs channel-first and channel-last permutations using the exact pattern from the issue report, and verifies a correct round-trip result.
- Fix misleading comment: remove 'then squeeze the result back' since concatenating [1]-shaped tensors along axis=0 already produces the correct 1-D vector of shape [N] with no squeezing needed. - Add non-empty guard (list_or_tuple and / elems_as_tensors and) in both code paths so that an empty input does not cause concat_v2 to be called with an empty list, which would raise an error.
concat_v2 requires at least 2 inputs (it raises: ValueError: List argument 'values' to 'ConcatV2' Op with length 1 shorter than minimum length 2.) When the Python list passed to tf.concat contains exactly one 0-D scalar tensor (the common case: [num_dims - 1]), the previous fix called concat_v2 with a single-element list and crashed. Fix: when there is only one scalar, return reshape(scalar, [1]) directly instead of wrapping it in concat_v2. Two or more scalars continue to use reshape+concat_v2 as before. Applied to both the eager fast-path and the graph-building path of _autopacking_helper.
The GPU CheckNumerics kernel emitted a LOG(ERROR) line on every NaN/Inf
detection, printing a host buffer address and raw indicator flags, e.g.
abnormal_detected_host @0x7fd6e3c05600 = {1, 0} <message>
immediately before setting an InvalidArgumentError that reports the same
condition in readable form. The log line therefore only duplicates the
error, floods stderr when anomalies are frequent, and exposes an address
that is of no use outside debugging this kernel.
Both siblings already report the condition without logging: the CPU
CheckNumerics kernel and CheckNumericsV2 set the status only. Demote the
line to VLOG(1) so it is silent by default and still reachable via
TF_CPP_VMODULE for kernel debugging.
Fixes #96524
Six lines of justification for a logging change was more than the code needs. Keep two lines explaining why this is intentionally not an error log, so the choice is not silently reverted, and drop the rest, which belongs in the PR rather than the source.
The testScalarConjugateTranspose test (regression test for issue #118345) was accidentally removed when inserting testConcatScalarRankDynamic above testScalar. Restore it to its original position between testScalar and _testError.
BincountReduceKernel, BincountColReduceKernel, and BincountColReduceSharedKernel only checked `bin < num_bins` before writing to the output buffer. A negative input value (e.g. from corrupt or adversarial input, or overflowed indices) satisfies that check and indexes before the start of the buffer, causing an out-of-bounds write on GPU. The CPU kernel already skips out-of-range bins in both directions, so this brings the GPU kernels in line with it by also checking `bin >= 0`. Fixes #103995
Per review feedback: cast bin and num_bins to their unsigned counterpart so the 0 <= bin < num_bins check is a single comparison instead of two, and avoid -Wtype-limits warnings if Tidx is ever unsigned.
Per review feedback, cover 64-bit index types in the BincountColReduceKernel/BincountColReduceSharedKernel negative-input regression tests, matching test_col_reduce_shared_memory / test_col_reduce_global_memory.
`SparseSegmentGradOpBase::Compute` sizes its output from `output_dim0`
(input 3, a runtime tensor) but takes its early return on
`N = indices.NumElements()`:
allocate_output(0, output_shape, &output); // [M, d1...] from output_dim0
if (M == 0 || N == 0) return; // output never written
M and N are independent, so an empty `indices` with `output_dim0 > 0`
returns a non-empty output that no code has written. `allocate_output`
does not zero the buffer, so the caller receives whatever the allocator
last held there.
The forward op in this same file already documents and handles exactly
this case, roughly 700 lines above:
// Note that we do not initialize the output buffer with a default value,
// so we need to explicitly set missing indices to the default value.
if (num_indices == 0) {
if (output_rows > 0) {
output->flat_outer_dims<T>().setConstant(default_value_);
}
return;
}
This mirrors that treatment for the gradient, where the identity is zero:
a row that received no contribution has a gradient of exactly zero.
Behaviour is unchanged for every defined input. When M == 0 the new guard
is false and the path is identical. When N > 0 the branch is not taken at
all, so the hot path is untouched. The only values that change are the
ones that were previously undefined.
Affects SparseSegmentSumGrad, SparseSegmentMeanGrad and
SparseSegmentSqrtNGrad, for float, double, half and bfloat16.
Not compile-tested (no Bazel environment) - please run CI.
Fixes host-side shape checking for ApplyAdadelta and ResourceApplyAdadelta.
… test Addresses both review comments: - The previous `output->flat_outer_dims<T>().setZero()` runs on the host CPU. `SparseSegmentGradOpBase` is templated on `Device` and is registered for GPU (`segment_reduction_ops_impl_5.cc`), so on GPU that wrote through a device pointer from the host. `functor::SetZeroFunctor<Device, T>` dispatches to the correct device. The guard is now `output->NumElements() > 0`, which is exact when the inner dimensions are empty as well. - Adds `testGradientEmptyIndicesWithNonZeroOutputDim` covering empty `indices` with `output_dim0 > 0` for sum/mean/sqrt_n over float16, bfloat16, float32 and float64. The test fails on an unpatched build, where the returned values are uninitialized heap contents rather than zero.
…d Conv2D When Grappler remapper fuses MatMul or Conv2D with Add, the bias tensor can have broadcast dimensions where leading axes are 1 (e.g. [1, N]). While mkl_matmul_op_fused allows bias_dims > 1 with leading 1s, InitBiasAddArgs in fused_eigen_output_kernels.h, GPU LaunchFusedMatMulOp, and LaunchFusedConv2DOp strictly required bias.dims() == 1. This caused InvalidArgumentError: bias must be 1-dimensional during graph execution in tf.function. Allow bias tensors with rank >= 1 whose leading dimensions are 1 in InitBiasAddArgs, GPU LaunchFusedMatMulOp, and LaunchFusedConv2DOp, aligning the Eigen and GPU kernels with the remapper and MKL implementations. Add test coverage in matmul_op_test.cc and matmul_op_test.py. Fixes #126661
…rflow write RaggedGather with int32 Tsplits could accumulate a total value count above int32, wrap it, allocate an undersized output, and then write past it (heap-buffer-overflow). Accumulate the count in int64 and reject counts that exceed the Tsplits maximum with InvalidArgumentError before allocating, matching the existing guard in ragged_range_op.cc. Adds a regression test that pins Tsplits=int32 via the raw op. Fixes #126135.
Address review feedback on #124017: The previous fix intercepted 0-D scalar tensors in the Python-level _autopacking_helper and replaced gen_array_ops.pack with reshape(scalar, [1]) + concat_v2. As pointed out in the review, this approach has several regressions: - Breaks static shape inference and constant propagation: PackOp retains known scalar values during shape inference, allowing TensorFlow to infer partial shapes (e.g. [None, 3]) for SparseTensor, RaggedTensor, and graph compilation. Replacing Pack with ConcatV2 loses this and produces [None, None] instead. - Incompatible with DT_RESOURCE and DT_VARIANT: ConcatV2 has no registered GPU kernel for these types, so tf.stack([r1, r2]) over resource handles fails with "No OpKernel was registered to support Op 'ConcatV2' with T=DT_RESOURCE". - Breaks graph pattern matching: TFLite OpHint AGGREGATE_STACK and XLA/Grappler passes pattern-match on Pack nodes; inserting Reshape + ConcatV2 prevents these transformations. Fix: Revert the Python-level _autopacking_helper changes entirely so that gen_array_ops.pack is always used, restoring full shape inference and type support. Add an explicit scalar-input branch in the C++ GPU PackOp kernel (tensorflow/core/kernels/pack_op.cc): when all inputs are 0-D scalars (first_input.dims() == 0, num > 1), write each scalar directly into the output flat vector using Eigen chip operations rather than routing through ConcatGPU, which was incompatible with 1-element-per-row configurations on some GPU drivers. The num == 1 scalar path already works via the existing CopyFrom reshape shortcut and is unchanged. Update the test docstring in testConcatScalarRankDynamic to reflect that the fix now lives in the C++ kernel rather than the Python autopacking helper.
The previous commit added a 0-D scalar shortcut in PackOp::Compute that
called:
output_vec.template chip<0>(i).device(c->eigen_device<Device>()) =
input.scalar<T>();
without restricting the code path to GPU-only. This caused compilation
failures on Linux x86-CPU, ARM64, and Windows x86 CI targets because:
1. The expression is instantiated for PackOp<CPUDevice, T> (where
CPUDevice is Eigen::ThreadPoolDevice). Eigen's device() assignment
asserts at compile time that the executor matches the device type; for
ThreadPoolDevice it fires:
error: static assertion failed: Default executor instantiated with
non-default device. You must #define EIGEN_USE_THREADS ...
2. Types like tensorflow::Variant and Eigen::QInt8 are only registered
for the CPU PackOp and never for GPU, so they are never affected by
the GPU bug; there is no reason to touch their code path at all.
Fix:
Wrap the entire scalar shortcut inside:
#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM
if constexpr (std::is_same_v<Device, GPUDevice>) { ... }
#endif
The outer preprocessor guard prevents the block from being compiled at
all on CPU-only builds. The inner if constexpr ensures that even in a
GPU-enabled build (where both PackOp<CPUDevice,...> and
PackOp<GPUDevice,...> are instantiated) the Eigen chip expression is
only compiled into the GPUDevice specialisation.
CPU builds and all non-scalar inputs are completely unaffected: they
fall through to the unchanged ConcatCPU / ConcatGPU generic path.
The previous fix placed the Eigen chip+device() assignment inside
pack_op.cc behind if constexpr (std::is_same_v<Device, GPUDevice>).
This still caused compilation failures on GPU builds because:
* pack_op.cc is a plain .cc file; NVCC does NOT define EIGEN_USE_GPU
for it, only for .cu.cc sources.
* Inside #if GOOGLE_CUDA, TF registers PackOp<CPUDevice, int32_t> for
the GPU device with HostMemory (line 208). That means the compiler
instantiates PackOp<CPUDevice,...>::Compute in this TU, and
even the discarded if-constexpr branch for GPUDevice causes Eigen to
emit its static assertion:
"Default executor instantiated with non-default device. You must
#define EIGEN_USE_THREADS, EIGEN_USE_GPU or EIGEN_USE_SYCL before
including Eigen headers."
Fix:
* Remove all Eigen device() expressions from pack_op.cc.
* Add pack_op_gpu.cu.cc which defines PackScalarsOnGPU<T>: a free
function that performs the scalar-element chip assignments using the
Eigen::GpuDevice. This file is compiled by NVCC/hipcc (as a .cu.cc
source), which defines EIGEN_USE_GPU before any Eigen headers are
included, satisfying Eigen's requirement.
* tf_kernel_library(name="pack_op", prefix="pack_op") in BUILD
automatically adds pack_op*.cu.cc to gpu_srcs, so no BUILD change
is needed.
* pack_op.cc forward-declares PackScalarsOnGPU<T> inside
#if GOOGLE_CUDA || TENSORFLOW_USE_ROCM and calls it when
std::is_same<Device, GPUDevice>::value && first_input.dims() == 0.
The regular std::is_same runtime check (not if constexpr) is used
here because no Eigen device expression appears in pack_op.cc any
more, so there is no risk of a static assertion.
* The CPU path (ConcatCPU) continues to handle all scalar inputs on
CPU unchanged.
* Explicit instantiations in pack_op_gpu.cu.cc cover every type
registered with REGISTER_GPU: TF_CALL_GPU_ALL_TYPES, int64, int16,
uint32, uint64, float8_e5m2, float8_e4m3fn.
Use 64-bit launch configuration with occupancy-aware GetGpuLaunchConfig64 to prevent integer overflow on large tensors in 2D and 3D dilation kernels, while staying within device resource limits.
The disclaimer paragraph in this file's header is word-scrambled and does not match the Apache 2.0 boilerplate every other file in the repository uses. It reads "To in writing unless required by applicable law or agreed, distributed on an, software distributed under the license is AS IS BASIS, WITHOUT OF ANY KIND WARRANTIES OR CONDITIONS ... the license you must see." Replace it with the standard wording and restore the blank line after the copyright line, so the header matches the neighbouring gpu_prim_helpers.h byte for byte from line 2 onward. The 2020 copyright year is kept. Searching the repository for each scrambled phrase returns only this file, and it has been this way since 25956c4 created it in January 2020, so nothing else needs the same treatment. Comments and code are untouched; this changes no behavior. Fixes #127660
PiperOrigin-RevId: 986917777
…ive-index PiperOrigin-RevId: 986921356
PiperOrigin-RevId: 986921845
…ar-rank-dynamic PiperOrigin-RevId: 986922239
PiperOrigin-RevId: 986922395
PiperOrigin-RevId: 986922513
…a module Imported from GitHub PR openxla/xla#49007 ### 📝 Summary of Changes `CudnnFusedConvRewriter::RunImpl` iterates over the non-fusion computations of a module and, for each of them, first runs `F8GraphConv`. If `F8GraphConv` rewrote anything, the function returned immediately, so the loop never reached the remaining computations. Any FP8 convolution in a later computation was left as a plain convolution: it got neither the `__cudnn$convForwardGraph` rewrite nor the bias/activation fusions, and the FP8 fallback later demoted it to a legacy non-FP8 `__cudnn$convForward` call. This change replaces the early `return` with `any_changed |= changed; continue;`. The `continue` keeps the existing behaviour for the rewritten computation (the remaining fusions only match the legacy conv custom calls and do not apply to graph convolutions), and the loop now processes every computation of the module. ### 🎯 Justification Hit in practice with JAX 0.10.2 on an RTX 5070 (sm_120, CUDA 12.9, cuDNN 9.24). A `jit` for MCTS self-play evaluates an FP8 convolutional network once in the entry computation and again inside a `lax.while_loop` body. Only one of the two copies was rewritten: the compiled HLO contained 20 `__cudnn$convForwardGraph` calls instead of 40, and the other 20 convolutions ran as legacy f32 convolutions. With this change all 40 are FP8 graph convolutions. Which copy loses depends on the computation order, so the workaround was to keep one copy out of FP8 entirely. Minimal JAX repro (same FP8 convolution in the entry computation and in a while body): ```python import jax, jax.numpy as jnp from jax import lax def conv(x, w): return lax.conv_general_dilated(x, w, (1, 1), "SAME", dimension_numbers=("NHWC", "HWIO", "NHWC"), preferred_element_type=jnp.float32) @jax.jit def f(x, w): y0 = conv(x, w) # entry computation def body(c): i, y = c return i + 1, y + conv(y.astype(jnp.float8_e4m3fn), w) # while body return lax.while_loop(lambda c: c[0] < 4, body, (0, y0))[1] x = jnp.ones((1, 6, 6, 128), jnp.float8_e4m3fn) w = jnp.ones((3, 3, 128, 128), jnp.float8_e4m3fn) hlo = f.lower(x, w).compile().as_text() print(hlo.count('custom_call_target="__cudnn$convForwardGraph"'), hlo.count('custom_call_target="__cudnn$convForward"')) ``` Before this change (jax-cuda12-pjrt 0.10.2): `1 1` (one graph convolution, one legacy convolution). With this change (plugin rebuilt from the same JAX release with the patched XLA): `2 0`. ### 🚀 Kind of Contribution 🐛 Bug Fix ### 🧪 Unit Tests `CudnnFusedConvRewriterTest.TestConvF8InMultipleComputations` in `cudnn_fused_conv_rewriter_test.cc`. The module has two `conditional` branches, each with an FP8 convolution on a different filter. After `ConvRewriter` and `CudnnFusedConvRewriter` (run with a Hopper compute capability, like the existing F8 tests on pre-FP8 hardware) both branch computations must contain a `__cudnn$convForwardGraph` custom call. Without the fix only the first branch visited is rewritten and the test fails on the second one. ### 🧪 Execution Tests On FP8-capable GPUs the same test additionally compiles the module through the full pipeline, checks with FileCheck that both convolutions are `__cudnn$convForwardGraph` and no legacy `__cudnn$convForward` remains, and compares the result against the interpreter (`RunAndCompare`, `ErrorSpec{0.15, 0.15}`). Tested locally on an RTX 5070 (sm_120) with the hermetic CUDA toolchain (cuDNN 9.12): `bazel test //xla/backends/gpu/transforms:cudnn_fused_conv_rewriter_test_nvgpu_any --test_filter='*TestConvF8InMultipleComputations*:*.TestConvF8:*TestConvScaledOutputF8'` passes (the new test runs the full-pipeline branch: `fp8_supported: true`). With the test but without the change in `cudnn_fused_conv_rewriter.cc` the new test fails as intended: `branch_b` keeps the legacy `__cudnn$convForward` custom call. Copybara import of the project: -- faf7a8a6ffa198a62772dcfe03ef90cfbb51eae3 by Yehor Cherednichenko <cherednichenko30igor@gmail.com>: [XLA:GPU] Rewrite FP8 convolutions in all computations of a module CudnnFusedConvRewriter::RunImpl returned as soon as F8GraphConv had rewritten the FP8 convolutions of one computation. FP8 convolutions in the remaining computations of the module (e.g. a while body next to the entry computation) were left unrewritten and later demoted to unfused f32 convolutions by the FP8 fallback. Continue with the next computation instead; the remaining fusions only match the legacy conv custom calls, so they still do not apply to the rewritten computation. Add TestConvF8InMultipleComputations: both branches of a conditional contain an FP8 convolution and both must get the ForwardGraph custom call. Merging this change closes #49007 PiperOrigin-RevId: 986927090
Imported from GitHub PR #123156 Fixes #103454 `MaxPoolWithArgmax` already supports `output_dtype=tf.int32` on CPU, but its gradient only registered an int64 argmax kernel. This caused automatic differentiation to fail with “No OpKernel was registered” when users selected int32 output indices. Template the CPU gradient implementation on the argmax dtype, register the existing CPU gradient for `Targmax=int32`, and extend the pooling gradient test matrix to validate both int32 and int64 CPU argmax tensors. GPU remains unchanged because its forward argmax kernel only supports int64. Validation: `git diff --check`; `python3 -m py_compile tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py`; `bazel build --config=macos_arm64 //tensorflow/core/kernels:pooling_ops` is running locally. Copybara import of the project: -- 7ce5452 by Saicharan Ramineni <csramineni@gmail.com>: Support int32 argmax in CPU max pool gradients Merging this change closes #123156 COPYBARA_INTEGRATE_REVIEW=#123156 from GodlyDonuts:fix/maxpool-int32-argmax 3ef9f72 PiperOrigin-RevId: 986934206
…empty PiperOrigin-RevId: 986939232
…-count-overflow PiperOrigin-RevId: 986953060
This unblocks the initial bulk migration, which will in turn unblock iterative migrations on a per-change basis. PiperOrigin-RevId: 986954805
…mmary PiperOrigin-RevId: 986956687
This was removed so it's failing on JAX builds. Just remove it for everything because we will likely bump soon and it provides little value PiperOrigin-RevId: 986984963
PiperOrigin-RevId: 987001752
PiperOrigin-RevId: 987004821
…ased on the fused computation Previously, `CseKey::AbslHashValue` in `HloCSE` fell into the `default:` branch for `HloOpcode::kFusion`, only hashing the top-level operands, output dimensions, and the fused computation's root opcode. As a result, distinct `kFusion` instructions sharing operands, output shape, and root opcode collided in `HloCSE`'s `representatives` hash table and triggered expensive `HloFusionInstruction::IdenticalSlowPath` (`HloComputation::Equal`) comparisons. This change factors out `CseKey::HashInstruction` and applies it to the internal instructions of `instruction->fused_instructions_computation()` (in post-order from `root_instruction()`) when hashing a `kFusion` instruction, skipping operand `unique_id()`s for internal instructions so equivalent `kFusion` instructions continue to hash identically while distinct fused computations are disambiguated. PiperOrigin-RevId: 987011851
…hlo_evaluator.h Imported from GitHub PR openxla/xla#49149 ## [XLA:Evaluator] Include `<cmath>` with `_USE_MATH_DEFINES` ### Problem `hlo_evaluator.h` defines `_USE_MATH_DEFINES` for MSVC but does not directly include `<cmath>`. This makes the availability of constants such as `M_PI` dependent on indirect include order. ### Solution Include `<cmath>` immediately after `_USE_MATH_DEFINES`, matching the existing pattern in `xla/fp_util.h`. ### Changes - `xla/hlo/evaluator/hlo_evaluator.h`: Add `#include <cmath>` after `_USE_MATH_DEFINES`. ### Testing - `git diff --check` — PASS - Unit tests not run locally; Bazel is unavailable. ### Risk Low risk. Header-only change with no API or runtime logic changes. Copybara import of the project: -- a8382c01fef2ee71002056b7bb91fe439175bcb3 by saketh reddy pingili <pingilisakethreddy@gmail.com>: [XLA:Evaluator] Include <cmath> with _USE_MATH_DEFINES in hlo_evaluator.h Merging this change closes #49149 PiperOrigin-RevId: 987040037
…fle-rounds PiperOrigin-RevId: 987044259
Add `tpu.vector_expand_load`, the logical counterpart of `tpu.vector_compress_store`. `result[i_0, i_1, ...]`, if not masked, is loaded from `base[j_0, j_1, ...]` where `j_n = indices_n + i_n` for `n != expand_dim`, and `j_n = indices_n + c` for `n == expand_dim`, with `c` the count of enabled mask values preceding position `i_n` in the slice. Elements at masked-off positions are undefined. PiperOrigin-RevId: 987058288
PiperOrigin-RevId: 987060773
…n barriers PiperOrigin-RevId: 987069321
PiperOrigin-RevId: 987077956
…tput-size PiperOrigin-RevId: 987096505
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 : )