From fea2878e2921ff48fd16df078fbc80ecd62ec518 Mon Sep 17 00:00:00 2001 From: Ayush Ojha Date: Sun, 1 Feb 2026 16:38:18 -0800 Subject: [PATCH 01/43] Validate scalar shape for input_min/input_max in QuantizeAndDequantizeV3 When axis=-1 and range_given=true, QuantizeAndDequantizeV3Op called .scalar() on input_min/input_max without first verifying they are scalars. This triggered a CHECK failure (process crash) instead of raising an InvalidArgumentError. The V2 op already had this validation. Fixes #99458 --- .../kernels/quantize_and_dequantize_op.cc | 10 ++++++++ .../quantization_ops/quantization_ops_test.py | 23 +++++++++++++++++++ 2 files changed, 33 insertions(+) diff --git a/tensorflow/core/kernels/quantize_and_dequantize_op.cc b/tensorflow/core/kernels/quantize_and_dequantize_op.cc index 64e7ec09c46eed..694b5c1f54cefa 100644 --- a/tensorflow/core/kernels/quantize_and_dequantize_op.cc +++ b/tensorflow/core/kernels/quantize_and_dequantize_op.cc @@ -296,6 +296,16 @@ class QuantizeAndDequantizeV3Op : public OpKernel { input_min_tensor = ctx->input(1); input_max_tensor = ctx->input(2); if (axis_ == -1) { + OP_REQUIRES( + ctx, TensorShapeUtils::IsScalar(input_min_tensor.shape()), + InvalidArgument("input_min must be a scalar when axis is not " + "specified. Got shape: ", + input_min_tensor.shape().DebugString())); + OP_REQUIRES( + ctx, TensorShapeUtils::IsScalar(input_max_tensor.shape()), + InvalidArgument("input_max must be a scalar when axis is not " + "specified. Got shape: ", + input_max_tensor.shape().DebugString())); const auto min_val = input_min_tensor.scalar()(); const auto max_val = input_max_tensor.scalar()(); OP_REQUIRES(ctx, min_val <= max_val, diff --git a/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py b/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py index 332d67ca76386b..846bf7dec6f778 100644 --- a/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py +++ b/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py @@ -487,6 +487,29 @@ def test_invalid_input_min_max_with_axis_specified(self): range_given=True, ) + @test_util.run_in_graph_and_eager_modes + def test_invalid_non_scalar_min_max_with_default_axis(self): + input_value = constant_op.constant( + [-0.8, -0.5, 0, 0.3, 0.8, -2.0], shape=(6,), dtype=dtypes.float32) + input_min = constant_op.constant( + [-127, -127], shape=(2,), dtype=dtypes.float32) + input_max = constant_op.constant( + [127, 127], shape=(2,), dtype=dtypes.float32) + num_bits = 8 + + with self.assertRaisesRegex( + (errors.InvalidArgumentError, ValueError), + "input_min must be a scalar"): + self.evaluate( + array_ops.quantize_and_dequantize_v3( + input_value, + input_min, + input_max, + num_bits=num_bits, + signed_input=True, + range_given=True, + )) + if __name__ == "__main__": googletest.main() From 7003ddecea2e5aa5d047190d6e50b61c7bad559d Mon Sep 17 00:00:00 2001 From: Adithyan AK Date: Tue, 7 Apr 2026 00:12:03 -0700 Subject: [PATCH 02/43] Fix OOB write in MaxPoolGradWithArgmax (CWE-787) The MaxPoolBackward GPU kernel at maxpooling_op_gpu.cu.cc:213 writes to `bottom_diff + offset + mask[index]` where `mask` comes from the user-supplied argmax tensor with no bounds checking. This allows an attacker-controlled OOB write via GpuAtomicAdd. The same file's MaxPoolGradBackward kernel (lines 332-350) already validates `read_index >= 0 && read_index < input_size` for the equivalent operation, confirming this was a missed check. On CPU, the equivalent code at maxpooling_op.cc:1090 uses CHECK() which kills the process with SIGABRT instead of returning an error. Fix: - GPU: Add `input_size` parameter to MaxPoolBackward kernel and validate `write_index >= 0 && write_index < input_size` before the GpuAtomicAdd, matching the existing GradGrad kernel pattern. - CPU: Replace CHECK() with a bounds check that skips invalid indices, preventing process crash on malformed argmax values. --- tensorflow/core/kernels/maxpooling_op.cc | 6 +++--- tensorflow/core/kernels/maxpooling_op_gpu.cu.cc | 10 +++++++--- 2 files changed, 10 insertions(+), 6 deletions(-) diff --git a/tensorflow/core/kernels/maxpooling_op.cc b/tensorflow/core/kernels/maxpooling_op.cc index a9de19492d1aff..1a129d008fc54d 100644 --- a/tensorflow/core/kernels/maxpooling_op.cc +++ b/tensorflow/core/kernels/maxpooling_op.cc @@ -1087,9 +1087,9 @@ struct LaunchMaxPoolingGradWithArgmax { const int64_t cur_batch = index / input_size_per_batch; grad_out_index += cur_batch * output_size_per_batch; } - CHECK(grad_out_index >= output_start && grad_out_index < output_end) - << "Invalid output gradient index: " << grad_out_index << ", " - << output_start << ", " << output_end; + if (grad_out_index < output_start || grad_out_index >= output_end) { + continue; + } grad_out_flat(grad_out_index) += grad_in_flat(index); } } diff --git a/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc b/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc index 99f52259e8e7a4..152ffcacfc254b 100644 --- a/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc +++ b/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc @@ -206,11 +206,15 @@ __global__ void MaxPoolBackward(const int nthreads, const int64_t* __restrict__ mask, const int top_offset, const int bottom_offset, dtype* __restrict__ bottom_diff, - const bool include_batch_in_index) { + const bool include_batch_in_index, + const int input_size) { GPU_1D_KERNEL_LOOP(index, nthreads) { const int offset = include_batch_in_index ? 0 : (index / top_offset) * bottom_offset; - GpuAtomicAdd(bottom_diff + offset + mask[index], top_diff[index]); + const int64_t write_index = offset + mask[index]; + if (write_index >= 0 && write_index < input_size) { + GpuAtomicAdd(bottom_diff + write_index, top_diff[index]); + } } } @@ -421,7 +425,7 @@ bool MaxPoolBackwardWithArgmax::operator()( MaxPoolBackward, (output_size + kThreadsPerBlock - 1) / kThreadsPerBlock, kThreadsPerBlock, 0, d.stream(), output_size, top_diff, mask, top_offset, bottom_offset, - bottom_diff, include_batch_in_index)); + bottom_diff, include_batch_in_index, input_size)); return d.ok(); } From 9ebce0d0fe2d72c6628d6e9955f668e9bbde7aa8 Mon Sep 17 00:00:00 2001 From: jmestwa-coder Date: Tue, 19 May 2026 18:45:14 +0530 Subject: [PATCH 03/43] Enforce runtime vector rank invariants in StringNGrams --- tensorflow/core/kernels/string_ngrams_op.cc | 8 ++++++++ tensorflow/python/ops/raw_ops_test.py | 18 ++++++++++++++++++ 2 files changed, 26 insertions(+) diff --git a/tensorflow/core/kernels/string_ngrams_op.cc b/tensorflow/core/kernels/string_ngrams_op.cc index 94d4009f84df07..f4d0765a378850 100644 --- a/tensorflow/core/kernels/string_ngrams_op.cc +++ b/tensorflow/core/kernels/string_ngrams_op.cc @@ -23,6 +23,7 @@ limitations under the License. #include "absl/strings/str_cat.h" #include "tensorflow/core/framework/op_kernel.h" #include "tensorflow/core/framework/op_requires.h" +#include "tensorflow/core/framework/tensor_shape.h" #include "tensorflow/core/platform/errors.h" #include "tensorflow/core/platform/types.h" @@ -73,10 +74,17 @@ class StringNGramsOp : public tensorflow::OpKernel { const tensorflow::Tensor* data; OP_REQUIRES_OK(context, context->input("data", &data)); + OP_REQUIRES(context, TensorShapeUtils::IsVector(data->shape()), + errors::InvalidArgument("data must be a vector, got shape: ", + data->shape().DebugString())); const auto& input_data = data->flat().data(); const tensorflow::Tensor* splits; OP_REQUIRES_OK(context, context->input("data_splits", &splits)); + OP_REQUIRES(context, TensorShapeUtils::IsVector(splits->shape()), + errors::InvalidArgument( + "data_splits must be a vector, got shape: ", + splits->shape().DebugString())); const auto& splits_vec = splits->flat(); // Validate that the splits are valid indices into data, only if there are diff --git a/tensorflow/python/ops/raw_ops_test.py b/tensorflow/python/ops/raw_ops_test.py index a6f19007b5d7e8..d5c17b0b22a15e 100644 --- a/tensorflow/python/ops/raw_ops_test.py +++ b/tensorflow/python/ops/raw_ops_test.py @@ -76,6 +76,24 @@ def testStringNGramsBadDataSplits(self, splits): pad_width=0, preserve_short_sequences=False)) + @parameterized.parameters( + ([[["aa"], ["bb"]]], [0, 2], "data must be a vector"), + (["aa", "bb"], [[0, 2]], "data_splits must be a vector"), + ) + def testStringNGramsRejectsNonVectorInputs( + self, data, data_splits, expected_error): + with self.assertRaisesRegex(errors.InvalidArgumentError, expected_error): + self.evaluate( + gen_string_ops.string_n_grams( + data=data, + data_splits=data_splits, + separator="", + ngram_widths=[1], + left_pad="", + right_pad="", + pad_width=0, + preserve_short_sequences=False)) + def testStringSplit(self): data = ["123456"] data_splits = [0, 1] From 7b7afc065e551d4b5285796794357274278606e3 Mon Sep 17 00:00:00 2001 From: Adithyan AK Date: Thu, 2 Jul 2026 16:29:55 -0700 Subject: [PATCH 04/43] Add regression test for MaxPoolGradWithArgmax OOB argmax indices Verifies that out-of-bounds argmax indices no longer crash the process on CPU (previously a CHECK failure / SIGABRT) or trigger silent OOB writes on GPU. Follow-up to the fix in 7003ddecea2 addressing dmiltr3's review request. --- .../kernel_tests/nn_ops/pooling_ops_test.py | 17 +++++++++++++++++ 1 file changed, 17 insertions(+) diff --git a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py index 781b41d6dbeb86..71fae40f175b76 100644 --- a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py @@ -2587,6 +2587,23 @@ def testMaxPoolGradWithArgmaxEagerShapeErrors(self): inp, grad, argmax, ksize=[1, 1, 1, 1], strides=[1, 1, 1, 1], padding="VALID") + def testMaxPoolGradWithArgmaxInvalidArgmaxIndices(self): + # Tests that out-of-bounds argmax indices do not cause process crash (on CPU) + # or OOB writes (on GPU). Let it skip invalid indices gracefully instead + # of hitting fatal CHECK failures. + with self.cached_session(): + inp = array_ops.ones((1, 3, 3, 1), dtype=dtypes.float32) + grad = array_ops.ones((1, 3, 3, 1), dtype=dtypes.float32) + # 999999 is out of bounds + argmax = constant_op.constant( + [999999] + [0] * 8, dtype=dtypes.int64, shape=[1, 3, 3, 1]) + + # This should execute without CRASH + res = gen_nn_ops.max_pool_grad_with_argmax( + inp, grad, argmax, ksize=[1, 1, 1, 1], strides=[1, 1, 1, 1], + padding="VALID") + self.evaluate(res) + def testAvgPoolGradInvalidInputShapeRaiseError(self): with self.assertRaises((ValueError, errors_impl.InvalidArgumentError)): with self.cached_session(): From 145da4b6b75c64377f390957d2be1284dbc94152 Mon Sep 17 00:00:00 2001 From: Ayush Ojha Date: Sun, 5 Jul 2026 01:25:17 -0700 Subject: [PATCH 05/43] test: accept graph quantize scalar shape error --- .../kernel_tests/quantization_ops/quantization_ops_test.py | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py b/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py index 846bf7dec6f778..841c4e908e2f32 100644 --- a/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py +++ b/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py @@ -499,7 +499,7 @@ def test_invalid_non_scalar_min_max_with_default_axis(self): with self.assertRaisesRegex( (errors.InvalidArgumentError, ValueError), - "input_min must be a scalar"): + "(input_min must be a scalar|Shape must be rank 0)"): self.evaluate( array_ops.quantize_and_dequantize_v3( input_value, From 832f9e1fdb8682bebe28d175cccfec4e5dc8acad Mon Sep 17 00:00:00 2001 From: Syed Mohammed Nayyar Date: Thu, 9 Jul 2026 23:45:43 +0530 Subject: [PATCH 06/43] handle graph-mode shape inference error in StringNGrams rank test --- tensorflow/python/ops/raw_ops_test.py | 34 ++++++++++++++++++--------- 1 file changed, 23 insertions(+), 11 deletions(-) diff --git a/tensorflow/python/ops/raw_ops_test.py b/tensorflow/python/ops/raw_ops_test.py index d5c17b0b22a15e..4211aedc6daf6d 100644 --- a/tensorflow/python/ops/raw_ops_test.py +++ b/tensorflow/python/ops/raw_ops_test.py @@ -82,17 +82,29 @@ def testStringNGramsBadDataSplits(self, splits): ) def testStringNGramsRejectsNonVectorInputs( self, data, data_splits, expected_error): - with self.assertRaisesRegex(errors.InvalidArgumentError, expected_error): - self.evaluate( - gen_string_ops.string_n_grams( - data=data, - data_splits=data_splits, - separator="", - ngram_widths=[1], - left_pad="", - right_pad="", - pad_width=0, - preserve_short_sequences=False)) + if context.executing_eagerly(): + with self.assertRaisesRegex(errors.InvalidArgumentError, expected_error): + self.evaluate( + gen_string_ops.string_n_grams( + data=data, + data_splits=data_splits, + separator="", + ngram_widths=[1], + left_pad="", + right_pad="", + pad_width=0, + preserve_short_sequences=False)) + else: + with self.assertRaisesRegex(ValueError, "Shape must be rank 1"): + gen_string_ops.string_n_grams( + data=data, + data_splits=data_splits, + separator="", + ngram_widths=[1], + left_pad="", + right_pad="", + pad_width=0, + preserve_short_sequences=False) def testStringSplit(self): data = ["123456"] From 32772f97940218b8a02445da7cac5c64a20ac34f Mon Sep 17 00:00:00 2001 From: mohammed adib Date: Fri, 17 Jul 2026 12:51:26 +0530 Subject: [PATCH 07/43] reject empty row_splits when decoding RaggedTensorVariant --- .../core/kernels/ragged_tensor_from_variant_op.cc | 6 ++++++ .../kernels/ragged_tensor_from_variant_op_test.cc | 11 +++++++++++ 2 files changed, 17 insertions(+) diff --git a/tensorflow/core/kernels/ragged_tensor_from_variant_op.cc b/tensorflow/core/kernels/ragged_tensor_from_variant_op.cc index e876e36fbc38ab..0accfb4db471e9 100644 --- a/tensorflow/core/kernels/ragged_tensor_from_variant_op.cc +++ b/tensorflow/core/kernels/ragged_tensor_from_variant_op.cc @@ -76,6 +76,12 @@ absl::Status RaggedComponentsFromVariant( "Ragged splits must have rank 1; encoded scalar element at index ", i, " has splits Tensor ", splits.DebugString())); } + if (splits.NumElements() < 1) { + return absl::InvalidArgumentError(absl::StrCat( + "Ragged splits must have at least one element; encoded scalar " + "element at index ", + i, " has splits Tensor ", splits.DebugString())); + } } } return absl::OkStatus(); diff --git a/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc b/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc index 3b9b7889e3e7d4..c0147f7f890e3a 100644 --- a/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc +++ b/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc @@ -548,6 +548,17 @@ TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitRankNotOne) { "Ragged splits must have rank 1")); } +TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitEmpty) { + RaggedTensorVariant encoded(Tensor(DT_INT32, {0}), {Tensor(DT_INT64, {0})}); + + int input_ragged_rank = 1; + int output_ragged_rank = 2; + BuildDecodeRaggedTensorGraph( + input_ragged_rank, output_ragged_rank, TensorShape({1}), {encoded}); + EXPECT_TRUE(absl::StartsWith(RunOpKernel().message(), + "Ragged splits must have at least one element")); +} + TEST_F(RaggedTensorFromVariantKernelTest, RaggedValuesTypeMismatch) { const std::vector component_split_1_1 = {0, 1}; const std::vector component_values_1 = {0}; From d7573bb7ded68408dc453cba58aeef8f263a947a Mon Sep 17 00:00:00 2001 From: mohammed adib Date: Mon, 20 Jul 2026 11:19:30 +0530 Subject: [PATCH 08/43] use CreateVariantFromRagged helper in RaggedSplitEmpty test --- .../core/kernels/ragged_tensor_from_variant_op_test.cc | 9 +++++++-- 1 file changed, 7 insertions(+), 2 deletions(-) diff --git a/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc b/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc index c0147f7f890e3a..fdf4da61fe6d50 100644 --- a/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc +++ b/tensorflow/core/kernels/ragged_tensor_from_variant_op_test.cc @@ -549,12 +549,17 @@ TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitRankNotOne) { } TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitEmpty) { - RaggedTensorVariant encoded(Tensor(DT_INT32, {0}), {Tensor(DT_INT64, {0})}); + const std::vector component_split_1_1 = {}; + const std::vector component_values_1 = {0}; + + auto variant_component_1 = CreateVariantFromRagged( + {component_split_1_1}, TensorShape({1}), component_values_1); int input_ragged_rank = 1; int output_ragged_rank = 2; BuildDecodeRaggedTensorGraph( - input_ragged_rank, output_ragged_rank, TensorShape({1}), {encoded}); + input_ragged_rank, output_ragged_rank, TensorShape({1}), + {variant_component_1}); EXPECT_TRUE(absl::StartsWith(RunOpKernel().message(), "Ragged splits must have at least one element")); } From 16924cf40f7ca7630c0837a6e11fdb7b3ac79e09 Mon Sep 17 00:00:00 2001 From: Yigtwxx Date: Thu, 16 Jul 2026 11:46:38 +0300 Subject: [PATCH 09/43] fix(scatter_nd): validate updates rank to avoid CHECK-fail crash ScatterNdOp::Compute and TensorScatterOp::Compute index updates.shape().dim_size(i) for every outer dimension of indices without first checking that updates has that many dimensions. When the rank of updates is smaller than the number of outer dimensions of indices (e.g. indices=[4,1,1], updates=[4]), TensorShape::dim_size hits CHECK(d < dims()) and aborts the process instead of returning an error. Add a rank guard before the loop so both ops return InvalidArgumentError, matching the validation already done in ValidateScatterNdUpdateShape on the deeper kernel path, and add regression tests. Fixes #93680 --- tensorflow/core/kernels/scatter_nd_op.cc | 16 +++++++++++ .../array_ops/scatter_nd_ops_test.py | 27 +++++++++++++++++++ 2 files changed, 43 insertions(+) diff --git a/tensorflow/core/kernels/scatter_nd_op.cc b/tensorflow/core/kernels/scatter_nd_op.cc index a1c8bdd66d15b8..9df981e8190268 100644 --- a/tensorflow/core/kernels/scatter_nd_op.cc +++ b/tensorflow/core/kernels/scatter_nd_op.cc @@ -142,6 +142,14 @@ class ScatterNdOp : public ScatterOpBase { const int64_t outer_dims = indices.shape().dims() - 1; + OP_REQUIRES(c, updates.shape().dims() >= outer_dims, + absl::InvalidArgumentError(absl::StrCat( + "Updates shape must have rank at least the number of " + "outer dimensions of indices (", outer_dims, + "). Found: updates shape=", + updates.shape().DebugString(), + ", indices shape=", indices.shape().DebugString()))); + for (int i = 0; i < outer_dims; ++i) { OP_REQUIRES( c, indices.shape().dim_size(i) == updates.shape().dim_size(i), @@ -216,6 +224,14 @@ class TensorScatterOp : public ScatterOpBase { const int64_t outer_dims = indices.shape().dims() - 1; + OP_REQUIRES(c, updates.shape().dims() >= outer_dims, + absl::InvalidArgumentError(absl::StrCat( + "Updates shape must have rank at least the number of " + "outer dimensions of indices (", outer_dims, + "). Found: updates shape=", + updates.shape().DebugString(), + ", indices shape=", indices.shape().DebugString()))); + for (int i = 0; i < outer_dims; ++i) { OP_REQUIRES(c, indices.shape().dim_size(i) == updates.shape().dim_size(i), absl::InvalidArgumentError(absl::StrCat( diff --git a/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py b/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py index 1c009421be177c..81fa6f352e6398 100644 --- a/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py +++ b/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py @@ -650,6 +650,19 @@ def testRank3InvalidShape2(self): r"Dimensions \[\d\,\d\) of input\[shape="): self.scatter_nd(indices, updates, shape) + def testUpdatesRankSmallerThanIndicesOuterDimsInvalid(self): + # Regression test for + # https://github.com/tensorflow/tensorflow/issues/93680: updates with + # rank smaller than the number of outer dimensions of indices used to + # crash with a CHECK failure instead of raising an error. + indices = array_ops.zeros([4, 1, 1], dtypes.int32) + updates = array_ops.zeros([4], dtypes.int32) + shape = np.array([8]) + with self.assertRaisesWithPredicateMatch( + (errors.InvalidArgumentError, ValueError), + r"rank at least|must match"): + self.scatter_nd(indices, updates, shape) + @parameterized.parameters(set((True, context.executing_eagerly()))) def testGradientsRank2ElementUpdate(self, use_tape): for dtype in GRADIENT_TESTS_DTYPES: @@ -830,6 +843,20 @@ class ScatterNdNonAliasingAddDeterminismTest(ScatterNdDeterminismTest, class ScatterNdTensorTest(test.TestCase): + @test_util.run_in_graph_and_eager_modes + def testUpdatesRankSmallerThanIndicesOuterDimsInvalid(self): + # Regression test for + # https://github.com/tensorflow/tensorflow/issues/93680: updates with + # rank smaller than the number of outer dimensions of indices used to + # crash with a CHECK failure instead of raising an error. + indices = constant_op.constant([[[4]], [[3]], [[1]], [[7]]]) + updates = constant_op.constant([9, 10, 11, 12]) + t = array_ops.ones([8], dtype=dtypes.int32) + with self.assertRaisesWithPredicateMatch( + (errors.InvalidArgumentError, ValueError), + r"rank at least|must match"): + self.evaluate(array_ops.tensor_scatter_update(t, indices, updates)) + @test_util.run_in_graph_and_eager_modes def testUpdateAddSub(self): for dtype in (dtypes.int32, dtypes.float32): From f22a66ee8feac202cae1f97eb73985f535c19022 Mon Sep 17 00:00:00 2001 From: Vishwak Thatikonda Date: Thu, 23 Jul 2026 11:56:06 -0700 Subject: [PATCH 10/43] Validate batch_index rank in Unbatch before accessing its dimensions UnbatchResource::Compute read batch_index_t.shape().dim_size(1) without checking that batch_index is a rank-2 matrix. A rank-1 batch_index aborted the process with a fatal CHECK failure (d < dims()) instead of raising a catchable InvalidArgumentError. A scalar batch_index took a different broken path: dim_size(0) returned an arbitrary garbage value, producing a nonsensical error message such as "Expected 0th dimension size to be no greater than 1; Got: 24040". Validate that batch_index is a matrix before any dim_size access. The existing testUnbatchInvalidIdArg passed a rank-3 batch_index that only incidentally survived the old checks; it now uses a valid rank-2 index so it still exercises the id validation it was written for. Fixes #104846 --- tensorflow/core/kernels/batch_kernels.cc | 8 ++++++++ tensorflow/python/ops/batch_ops_test.py | 22 +++++++++++++++++++++- 2 files changed, 29 insertions(+), 1 deletion(-) diff --git a/tensorflow/core/kernels/batch_kernels.cc b/tensorflow/core/kernels/batch_kernels.cc index d6ee848aa7d017..f900c764139d77 100644 --- a/tensorflow/core/kernels/batch_kernels.cc +++ b/tensorflow/core/kernels/batch_kernels.cc @@ -828,6 +828,14 @@ class UnbatchResource : public ResourceBase { const Tensor& data_t = context->input(0); const Tensor& batch_index_t = context->input(1); + // The rank must be validated before any dim_size access, which has + // undefined behavior for out-of-range dimension indices. + if (!TensorShapeUtils::IsMatrix(batch_index_t.shape())) { + return absl::InvalidArgumentError( + absl::StrCat("Wrong shape for index tensor. Expected a matrix of " + "shape [batch_size, 3]; Got: ", + batch_index_t.shape().DebugString(), ".")); + } if (batch_index_t.shape().dim_size(0) > data_t.shape().dim_size(0)) { return absl::InvalidArgumentError(absl::StrCat( "Wrong shape for index tensor. Expected 0th dimension size to be no " diff --git a/tensorflow/python/ops/batch_ops_test.py b/tensorflow/python/ops/batch_ops_test.py index 15a1a71a93bb63..4c09c06d13b797 100644 --- a/tensorflow/python/ops/batch_ops_test.py +++ b/tensorflow/python/ops/batch_ops_test.py @@ -242,7 +242,7 @@ def testUnbatchInvalidIdArg(self): batched_tensor = constant_op.constant( value=np.random.random(size=(3, 3, 1)), dtype=dtypes.float64) batched_index = constant_op.constant( - value=np.random.randint(0, 100, size=(3, 3, 1)), dtype=dtypes.int64) + value=np.random.randint(0, 100, size=(3, 3)), dtype=dtypes.int64) arg_id = constant_op.constant( value=np.random.randint(0, 100, size=(3, 3, 1)), dtype=dtypes.int64) @@ -256,6 +256,26 @@ def testUnbatchInvalidIdArg(self): container="", shared_name="") + def testUnbatchInvalidIndexRank(self): + # Regression test for GitHub issue 104846: a batch_index tensor that is + # not a rank-2 matrix used to crash the process with a fatal CHECK + # failure instead of raising InvalidArgumentError. + if context.executing_eagerly(): + for bad_index in ( + constant_op.constant([0], dtype=dtypes.int64), + constant_op.constant(0, dtype=dtypes.int64), + ): + with self.assertRaisesRegex( + errors.InvalidArgumentError, + r"Expected a matrix of shape \[batch_size, 3\]"): + batch_ops.unbatch( + batched_tensor=constant_op.constant([1], dtype=dtypes.int32), + batch_index=bad_index, + id=constant_op.constant(0, dtype=dtypes.int64), + timeout_micros=0, + container="", + shared_name="") + def testBatchDecoratedWithCapturedInput(self): """Tests that the batch_function decorator works.""" if context.executing_eagerly(): From 448305d33bb26ae7f3511334c3b7e639899662d8 Mon Sep 17 00:00:00 2001 From: Yigtwxx Date: Sat, 25 Jul 2026 00:12:45 +0300 Subject: [PATCH 11/43] test(scatter_nd): address review feedback on rank-validation tests The internal presubmits run these regression tests under XLA compilation, where tf2xla lowers the op itself and reports the shape mismatch with a different message ("Must have updates.shape = ...") than the CPU/GPU kernel ("rank at least ...") or graph-mode shape inference ("must match"). Accept all three so the assertion holds on every execution path. Also run the tf.scatter_nd regression test in graph mode as well as eager, matching the tensor_scatter_update one, so graph-mode shape inference is covered too. --- .../kernel_tests/array_ops/scatter_nd_ops_test.py | 11 +++++++++-- 1 file changed, 9 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py b/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py index 81fa6f352e6398..302151dc362fb0 100644 --- a/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py +++ b/tensorflow/python/kernel_tests/array_ops/scatter_nd_ops_test.py @@ -650,6 +650,7 @@ def testRank3InvalidShape2(self): r"Dimensions \[\d\,\d\) of input\[shape="): self.scatter_nd(indices, updates, shape) + @test_util.run_in_graph_and_eager_modes def testUpdatesRankSmallerThanIndicesOuterDimsInvalid(self): # Regression test for # https://github.com/tensorflow/tensorflow/issues/93680: updates with @@ -658,9 +659,12 @@ def testUpdatesRankSmallerThanIndicesOuterDimsInvalid(self): indices = array_ops.zeros([4, 1, 1], dtypes.int32) updates = array_ops.zeros([4], dtypes.int32) shape = np.array([8]) + # The message differs per path: the CPU/GPU kernel reports "rank at + # least", graph-mode shape inference reports "must match", and the + # tf2xla lowering reports "Must have updates.shape = ...". with self.assertRaisesWithPredicateMatch( (errors.InvalidArgumentError, ValueError), - r"rank at least|must match"): + r"rank at least|must match|Must have updates\.shape"): self.scatter_nd(indices, updates, shape) @parameterized.parameters(set((True, context.executing_eagerly()))) @@ -852,9 +856,12 @@ def testUpdatesRankSmallerThanIndicesOuterDimsInvalid(self): indices = constant_op.constant([[[4]], [[3]], [[1]], [[7]]]) updates = constant_op.constant([9, 10, 11, 12]) t = array_ops.ones([8], dtype=dtypes.int32) + # The message differs per path: the CPU/GPU kernel reports "rank at + # least", graph-mode shape inference reports "must match", and the + # tf2xla lowering reports "Must have updates.shape = ...". with self.assertRaisesWithPredicateMatch( (errors.InvalidArgumentError, ValueError), - r"rank at least|must match"): + r"rank at least|must match|Must have updates\.shape"): self.evaluate(array_ops.tensor_scatter_update(t, indices, updates)) @test_util.run_in_graph_and_eager_modes From efadb4b9b559c7896d3ef0139aefa0b55acc2d19 Mon Sep 17 00:00:00 2001 From: Vishwak Thatikonda Date: Thu, 13 Aug 2026 11:44:22 -0700 Subject: [PATCH 12/43] Validate reverse axes for scalar and empty inputs ReverseV2 returned the input unchanged for scalar and empty tensors before looking at the axes, so out-of-range axes went unreported in eager execution while shape inference rejected the same call in graph mode. A scalar has no valid axis at all, yet tf.reverse(scalar, axis=[1,2,3]) returned the scalar; an empty tensor with an out-of-range axis behaved the same way. Validate the axes before taking the shortcut for inputs that have nothing to reverse, then return early only for the data movement. The checks and their messages are unchanged, so eager execution now reports what graph mode already did. An empty axis list stays valid for every input, and in-range axes on empty tensors keep working. Most of the diff is re-indentation from removing the enclosing else block. Fixes #110038 --- tensorflow/core/kernels/reverse_op.cc | 93 ++++++++++--------- .../kernel_tests/array_ops/array_ops_test.py | 35 +++++++ 2 files changed, 85 insertions(+), 43 deletions(-) diff --git a/tensorflow/core/kernels/reverse_op.cc b/tensorflow/core/kernels/reverse_op.cc index 87bf875cd45e22..89e021601cbcf7 100644 --- a/tensorflow/core/kernels/reverse_op.cc +++ b/tensorflow/core/kernels/reverse_op.cc @@ -249,61 +249,68 @@ class ReverseV2Op : public OpKernel { const Tensor& input = context->input(0); const Tensor& sparse_dims = context->input(1); + const int input_dims = input.dims(); + const TensorShape& sparse_dims_shape = sparse_dims.shape(); + const auto& axes_sparse_flat = sparse_dims.flat(); + + // The axes are validated before the shortcut for scalar and empty inputs + // below. Those inputs have nothing to reverse, but their axes can still be + // out of range, and shape inference already rejects them in graph mode. + // Validating here keeps eager execution consistent with that. Note that a + // scalar has no valid axis at all, so any axis is rejected for it, while + // an empty axis list stays valid for every input. + OP_REQUIRES(context, TensorShapeUtils::IsVector(sparse_dims_shape), + absl::InvalidArgumentError(absl::StrCat( + "'dims' must be 1-dimension, not ", sparse_dims.dims()))); + absl::InlinedVector axes_dense(input_dims, false); + for (int dummy = 0; dummy < axes_sparse_flat.size(); dummy++) { + Tidx axis = internal::SubtleMustCopy(axes_sparse_flat(dummy)); + Tidx canonical_axis = axis < 0 ? input_dims + axis : axis; + OP_REQUIRES(context, canonical_axis >= 0 && canonical_axis < input_dims, + errors::InvalidArgument("'axis'[", dummy, "] = ", axis, + " is out of valid range [", 0, ", ", + input_dims - 1)); + OP_REQUIRES(context, !axes_dense[canonical_axis], + errors::InvalidArgument("axis ", canonical_axis, + " specified more than once.")); + axes_dense[canonical_axis] = true; + } + if (TensorShapeUtils::IsScalar(input.shape()) || input.NumElements() == 0) { context->set_output(0, input); - } else { - const int input_dims = input.dims(); - const TensorShape& sparse_dims_shape = sparse_dims.shape(); - const auto& axes_sparse_flat = sparse_dims.flat(); - - OP_REQUIRES(context, TensorShapeUtils::IsVector(sparse_dims_shape), - absl::InvalidArgumentError(absl::StrCat( - "'dims' must be 1-dimension, not ", sparse_dims.dims()))); - absl::InlinedVector axes_dense(input_dims, false); - for (int dummy = 0; dummy < axes_sparse_flat.size(); dummy++) { - Tidx axis = internal::SubtleMustCopy(axes_sparse_flat(dummy)); - Tidx canonical_axis = axis < 0 ? input_dims + axis : axis; - OP_REQUIRES(context, canonical_axis >= 0 && canonical_axis < input_dims, - errors::InvalidArgument("'axis'[", dummy, "] = ", axis, - " is out of valid range [", 0, ", ", - input_dims - 1)); - OP_REQUIRES(context, !axes_dense[canonical_axis], - errors::InvalidArgument("axis ", canonical_axis, - " specified more than once.")); - axes_dense[canonical_axis] = true; - } + return; + } - OP_REQUIRES(context, input_dims <= 8, - absl::UnimplementedError( - "reverse is not implemented for tensors of rank > 8.")); + OP_REQUIRES(context, input_dims <= 8, + absl::UnimplementedError( + "reverse is not implemented for tensors of rank > 8.")); - Tensor* output = nullptr; - OP_REQUIRES_OK(context, - context->allocate_output(0, input.shape(), &output)); + Tensor* output = nullptr; + OP_REQUIRES_OK(context, + context->allocate_output(0, input.shape(), &output)); - // TODO(cwhipkey): we can do dimension folding to reduce, e.g., a reverse - // of a single dimension to the dims=3 or dims=2 case, regardless of the - // number of dimensions in the tensor. This would let some ops use faster - // lower-dimension code (and use optimized versions). + // TODO(cwhipkey): we can do dimension folding to reduce, e.g., a reverse + // of a single dimension to the dims=3 or dims=2 case, regardless of the + // number of dimensions in the tensor. This would let some ops use faster + // lower-dimension code (and use optimized versions). #define HANDLE_REVERSE(NDIMS) \ case NDIMS: \ HandleReverseV2Case(context, axes_dense, output); \ return; - switch (input_dims) { - HANDLE_REVERSE(0); - HANDLE_REVERSE(1); - HANDLE_REVERSE(2); - HANDLE_REVERSE(3); - HANDLE_REVERSE(4); - HANDLE_REVERSE(5); - HANDLE_REVERSE(6); - HANDLE_REVERSE(7); - HANDLE_REVERSE(8); - } -#undef HANDLE_REVERSE + switch (input_dims) { + HANDLE_REVERSE(0); + HANDLE_REVERSE(1); + HANDLE_REVERSE(2); + HANDLE_REVERSE(3); + HANDLE_REVERSE(4); + HANDLE_REVERSE(5); + HANDLE_REVERSE(6); + HANDLE_REVERSE(7); + HANDLE_REVERSE(8); } +#undef HANDLE_REVERSE } }; diff --git a/tensorflow/python/kernel_tests/array_ops/array_ops_test.py b/tensorflow/python/kernel_tests/array_ops/array_ops_test.py index 8a25816f0d493b..34cdee33daaf96 100644 --- a/tensorflow/python/kernel_tests/array_ops/array_ops_test.py +++ b/tensorflow/python/kernel_tests/array_ops/array_ops_test.py @@ -591,6 +591,41 @@ def testReverseInvalidShape(self): v = array_ops.reverse_v2(x, axis=[1]) self.assertAllEqual(self.evaluate(v), v) + def testReverseScalarOutOfRangeAxis(self): + # Regression test for GitHub issue 110038: a scalar has no valid axis, so + # any axis must be rejected. The kernel returned the input unchanged + # instead, which disagreed with the error shape inference raises for the + # same call in graph mode. + for axis in ([0], [1, 2, 3], [-1]): + with self.subTest(axis=axis): + with self.assertRaisesRegex( + (ValueError, errors.InvalidArgumentError), "out of valid range"): + self.evaluate( + array_ops.reverse_v2(constant_op.constant(4.0), axis=axis)) + + def testReverseScalarEmptyAxisIsValid(self): + # An empty axis list reverses nothing and stays valid for any input. + x = constant_op.constant(4.0) + self.assertAllEqual(4.0, self.evaluate(array_ops.reverse_v2(x, axis=[]))) + + def testReverseEmptyTensorOutOfRangeAxis(self): + # The same validation gap applied to empty tensors, whose axes can be out + # of range even though there is nothing to reverse. + for shape, axis in (([0], [5]), ([0, 3], [7]), ([0, 3], [-4])): + with self.subTest(shape=shape, axis=axis): + x = array_ops.zeros(shape, dtype=dtypes.float32) + with self.assertRaisesRegex( + (ValueError, errors.InvalidArgumentError), "out of valid range"): + self.evaluate(array_ops.reverse_v2(x, axis=axis)) + + def testReverseEmptyTensorValidAxis(self): + # In-range axes on empty tensors keep working. + for shape, axis in (([0], [0]), ([0, 3], [1]), ([0, 3], [-1])): + with self.subTest(shape=shape, axis=axis): + x = array_ops.zeros(shape, dtype=dtypes.float32) + self.assertAllEqual( + np.zeros(shape), self.evaluate(array_ops.reverse_v2(x, axis=axis))) + class MeshgridTest(test_util.TensorFlowTestCase): From 86d7965e79954410bee1ed44f2345f60ea996a60 Mon Sep 17 00:00:00 2001 From: Veer Jain Date: Sat, 15 Aug 2026 12:15:34 -0700 Subject: [PATCH 13/43] fix(rnn): bounds check seq_len_max in BlockLSTM ops to prevent OOB access --- tensorflow/core/kernels/rnn/lstm_ops.cc | 10 ++++++++++ 1 file changed, 10 insertions(+) diff --git a/tensorflow/core/kernels/rnn/lstm_ops.cc b/tensorflow/core/kernels/rnn/lstm_ops.cc index bc2398f7b84ab3..60f471050a8954 100644 --- a/tensorflow/core/kernels/rnn/lstm_ops.cc +++ b/tensorflow/core/kernels/rnn/lstm_ops.cc @@ -1058,6 +1058,11 @@ class BlockLSTMOp : public OpKernel { const Device& device = ctx->eigen_device(); const int64_t seq_len_max = seq_len_max_tensor->scalar()(); + OP_REQUIRES( + ctx, seq_len_max >= 0 && seq_len_max <= timelen, + absl::InvalidArgumentError(absl::StrCat( + "seq_len_max must be between 0 and ", timelen, " but is ", + seq_len_max))); SliceHelper slicer(ctx); for (int64_t t = 0; t < seq_len_max; ++t) { const Tensor x_tensor = slicer.InputSlice(*x, t, "x"); @@ -1377,6 +1382,11 @@ class BlockLSTMGradOp : public OpKernel { functor::TensorZero()(device, b_grad_tensor->flat()); const int64_t seq_len_max = seq_len_max_tensor->scalar()(); + OP_REQUIRES( + ctx, seq_len_max >= 0 && seq_len_max <= timelen, + absl::InvalidArgumentError(absl::StrCat( + "seq_len_max must be between 0 and ", timelen, " but is ", + seq_len_max))); SliceHelper slicer(ctx); for (int64_t t = seq_len_max - 1; t >= 0; --t) { const Tensor& x_tensor = slicer.InputSlice(*x, t, "x"); From c841dacd0a05a0f10deefbd1bded2a5513ca0f17 Mon Sep 17 00:00:00 2001 From: Veer Jain Date: Sun, 23 Aug 2026 11:32:29 -0700 Subject: [PATCH 14/43] test: add bounds validation tests for BlockLSTM seq_len_max --- .../kernel_tests/nn_ops/rnn_cell_test.py | 131 ++++++++++++++++++ 1 file changed, 131 insertions(+) diff --git a/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py b/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py index a1f4d741ba2782..32fd04c9e5d31d 100644 --- a/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py @@ -1468,7 +1468,138 @@ def testLSTMBlockInvalidArgument(self): ) ) + @test_util.run_in_graph_and_eager_modes + def testBlockLSTMSeqLenMaxBounds(self): + timelen, batch_size, input_size, cell_size = 4, 2, 3, 5 + x = constant_op.constant(0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) + cs_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + h_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + w = constant_op.constant(0.1, shape=[input_size + cell_size, 4 * cell_size], dtype=dtypes.float32) + wci = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + wcf = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + wco = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + b = constant_op.constant(0.1, shape=[4 * cell_size], dtype=dtypes.float32) + + # Valid boundary cases: seq_len_max == 0 and seq_len_max == timelen + for valid_seq_len_max in [0, 2, timelen]: + res = self.evaluate( + gen_rnn_ops.BlockLSTM( + seq_len_max=constant_op.constant(valid_seq_len_max, dtype=dtypes.int64), + x=x, + cs_prev=cs_prev, + h_prev=h_prev, + w=w, + wci=wci, + wcf=wcf, + wco=wco, + b=b, + forget_bias=0.0, + cell_clip=-1.0, + use_peephole=False, + ) + ) + self.assertEqual(len(res), 7) + + # Invalid cases: seq_len_max < 0 or seq_len_max > timelen + for invalid_seq_len_max in [-1, timelen + 1, 100]: + with self.assertRaisesRegex( + (ValueError, errors_impl.InvalidArgumentError), + r"seq_len_max must be between 0 and", + ): + self.evaluate( + gen_rnn_ops.BlockLSTM( + seq_len_max=constant_op.constant(invalid_seq_len_max, dtype=dtypes.int64), + x=x, + cs_prev=cs_prev, + h_prev=h_prev, + w=w, + wci=wci, + wcf=wcf, + wco=wco, + b=b, + forget_bias=0.0, + cell_clip=-1.0, + use_peephole=False, + ) + ) + @test_util.run_in_graph_and_eager_modes + def testBlockLSTMGradSeqLenMaxBounds(self): + timelen, batch_size, input_size, cell_size = 4, 2, 3, 5 + x = constant_op.constant(0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) + cs_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + h_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + w = constant_op.constant(0.1, shape=[input_size + cell_size, 4 * cell_size], dtype=dtypes.float32) + wci = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + wcf = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + wco = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) + b = constant_op.constant(0.1, shape=[4 * cell_size], dtype=dtypes.float32) + i = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + cs = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + f = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + o = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + ci = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + co = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + h = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + cs_grad = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + h_grad = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + + # Valid boundary cases: seq_len_max == 0 and seq_len_max == timelen + for valid_seq_len_max in [0, 2, timelen]: + res = self.evaluate( + gen_rnn_ops.BlockLSTMGrad( + seq_len_max=constant_op.constant(valid_seq_len_max, dtype=dtypes.int64), + x=x, + cs_prev=cs_prev, + h_prev=h_prev, + w=w, + wci=wci, + wcf=wcf, + wco=wco, + b=b, + i=i, + cs=cs, + f=f, + o=o, + ci=ci, + co=co, + h=h, + cs_grad=cs_grad, + h_grad=h_grad, + use_peephole=False, + ) + ) + self.assertEqual(len(res), 8) + + # Invalid cases: seq_len_max < 0 or seq_len_max > timelen + for invalid_seq_len_max in [-1, timelen + 1, 100]: + with self.assertRaisesRegex( + (ValueError, errors_impl.InvalidArgumentError), + r"seq_len_max must be between 0 and", + ): + self.evaluate( + gen_rnn_ops.BlockLSTMGrad( + seq_len_max=constant_op.constant(invalid_seq_len_max, dtype=dtypes.int64), + x=x, + cs_prev=cs_prev, + h_prev=h_prev, + w=w, + wci=wci, + wcf=wcf, + wco=wco, + b=b, + i=i, + cs=cs, + f=f, + o=o, + ci=ci, + co=co, + h=h, + cs_grad=cs_grad, + h_grad=h_grad, + use_peephole=False, + ) + ) class BidirectionalRNNTest(test.TestCase): def setUp(self): From f8fbb9611a80a12b015f4ac9f860026b8c827d77 Mon Sep 17 00:00:00 2001 From: Veer Jain Date: Sun, 23 Aug 2026 14:29:13 -0700 Subject: [PATCH 15/43] fix(test): wrap long lines in rnn_cell_test.py to pass PyLint --- .../kernel_tests/nn_ops/rnn_cell_test.py | 71 +++++++++++++------ 1 file changed, 48 insertions(+), 23 deletions(-) diff --git a/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py b/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py index 32fd04c9e5d31d..b60baefd17e977 100644 --- a/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py @@ -1471,20 +1471,27 @@ def testLSTMBlockInvalidArgument(self): @test_util.run_in_graph_and_eager_modes def testBlockLSTMSeqLenMaxBounds(self): timelen, batch_size, input_size, cell_size = 4, 2, 3, 5 - x = constant_op.constant(0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) - cs_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) - h_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) - w = constant_op.constant(0.1, shape=[input_size + cell_size, 4 * cell_size], dtype=dtypes.float32) + x = constant_op.constant( + 0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) + cs_prev = constant_op.constant( + 0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + h_prev = constant_op.constant( + 0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + w = constant_op.constant( + 0.1, shape=[input_size + cell_size, 4 * cell_size], + dtype=dtypes.float32) wci = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) wcf = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) wco = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) - b = constant_op.constant(0.1, shape=[4 * cell_size], dtype=dtypes.float32) + b = constant_op.constant( + 0.1, shape=[4 * cell_size], dtype=dtypes.float32) # Valid boundary cases: seq_len_max == 0 and seq_len_max == timelen for valid_seq_len_max in [0, 2, timelen]: res = self.evaluate( gen_rnn_ops.BlockLSTM( - seq_len_max=constant_op.constant(valid_seq_len_max, dtype=dtypes.int64), + seq_len_max=constant_op.constant( + valid_seq_len_max, dtype=dtypes.int64), x=x, cs_prev=cs_prev, h_prev=h_prev, @@ -1508,7 +1515,8 @@ def testBlockLSTMSeqLenMaxBounds(self): ): self.evaluate( gen_rnn_ops.BlockLSTM( - seq_len_max=constant_op.constant(invalid_seq_len_max, dtype=dtypes.int64), + seq_len_max=constant_op.constant( + invalid_seq_len_max, dtype=dtypes.int64), x=x, cs_prev=cs_prev, h_prev=h_prev, @@ -1526,29 +1534,45 @@ def testBlockLSTMSeqLenMaxBounds(self): @test_util.run_in_graph_and_eager_modes def testBlockLSTMGradSeqLenMaxBounds(self): timelen, batch_size, input_size, cell_size = 4, 2, 3, 5 - x = constant_op.constant(0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) - cs_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) - h_prev = constant_op.constant(0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) - w = constant_op.constant(0.1, shape=[input_size + cell_size, 4 * cell_size], dtype=dtypes.float32) + x = constant_op.constant( + 0.1, shape=[timelen, batch_size, input_size], dtype=dtypes.float32) + cs_prev = constant_op.constant( + 0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + h_prev = constant_op.constant( + 0.1, shape=[batch_size, cell_size], dtype=dtypes.float32) + w = constant_op.constant( + 0.1, shape=[input_size + cell_size, 4 * cell_size], + dtype=dtypes.float32) wci = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) wcf = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) wco = constant_op.constant(0.1, shape=[cell_size], dtype=dtypes.float32) - b = constant_op.constant(0.1, shape=[4 * cell_size], dtype=dtypes.float32) - i = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - cs = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - f = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - o = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - ci = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - co = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - h = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - cs_grad = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) - h_grad = constant_op.constant(0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + b = constant_op.constant( + 0.1, shape=[4 * cell_size], dtype=dtypes.float32) + i = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + cs = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + f = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + o = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + ci = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + co = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + h = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + cs_grad = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) + h_grad = constant_op.constant( + 0.1, shape=[timelen, batch_size, cell_size], dtype=dtypes.float32) # Valid boundary cases: seq_len_max == 0 and seq_len_max == timelen for valid_seq_len_max in [0, 2, timelen]: res = self.evaluate( gen_rnn_ops.BlockLSTMGrad( - seq_len_max=constant_op.constant(valid_seq_len_max, dtype=dtypes.int64), + seq_len_max=constant_op.constant( + valid_seq_len_max, dtype=dtypes.int64), x=x, cs_prev=cs_prev, h_prev=h_prev, @@ -1579,7 +1603,8 @@ def testBlockLSTMGradSeqLenMaxBounds(self): ): self.evaluate( gen_rnn_ops.BlockLSTMGrad( - seq_len_max=constant_op.constant(invalid_seq_len_max, dtype=dtypes.int64), + seq_len_max=constant_op.constant( + invalid_seq_len_max, dtype=dtypes.int64), x=x, cs_prev=cs_prev, h_prev=h_prev, From 40ad276180ad28eea27166be1daaf6fe6e4b2a65 Mon Sep 17 00:00:00 2001 From: Veer Jain Date: Mon, 24 Aug 2026 14:27:23 -0400 Subject: [PATCH 16/43] fix(test): update expected error message for seq_len_max bounds check --- tensorflow/python/ops/rnn_grad_test.py | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/tensorflow/python/ops/rnn_grad_test.py b/tensorflow/python/ops/rnn_grad_test.py index 59ed5dee3970e3..5f7f068aaaa6cf 100644 --- a/tensorflow/python/ops/rnn_grad_test.py +++ b/tensorflow/python/ops/rnn_grad_test.py @@ -120,7 +120,7 @@ def testLSTMBlockCell(self): def testBlockLSTMSeqLenMaxTooLarge(self): w, b, x, cs_prev, h_prev, w_peephole = self._block_lstm_inputs() with self.assertRaisesRegex( - errors_impl.InvalidArgumentError, "seq_len_max must be <= timelen" + errors_impl.InvalidArgumentError, r"seq_len_max must be between 0 and" ): self.evaluate( self._block_lstm(w, b, x, cs_prev, h_prev, w_peephole, seq_len_max=10) @@ -129,7 +129,7 @@ def testBlockLSTMSeqLenMaxTooLarge(self): def testBlockLSTMSeqLenMaxNegative(self): w, b, x, cs_prev, h_prev, w_peephole = self._block_lstm_inputs() with self.assertRaisesRegex( - errors_impl.InvalidArgumentError, "seq_len_max must be >= 0" + errors_impl.InvalidArgumentError, r"seq_len_max must be between 0 and" ): self.evaluate( self._block_lstm(w, b, x, cs_prev, h_prev, w_peephole, seq_len_max=-1) @@ -137,13 +137,13 @@ def testBlockLSTMSeqLenMaxNegative(self): def testBlockLSTMGradSeqLenMaxTooLarge(self): with self.assertRaisesRegex( - errors_impl.InvalidArgumentError, "seq_len_max must be <= timelen" + errors_impl.InvalidArgumentError, r"seq_len_max must be between 0 and" ): self.evaluate(self._block_lstm_grad(seq_len_max=10)) def testBlockLSTMGradSeqLenMaxNegative(self): with self.assertRaisesRegex( - errors_impl.InvalidArgumentError, "seq_len_max must be >= 0" + errors_impl.InvalidArgumentError, r"seq_len_max must be between 0 and" ): self.evaluate(self._block_lstm_grad(seq_len_max=-1)) From 67fa3b04e27c8ff778124e623aa77eeb0cdca290 Mon Sep 17 00:00:00 2001 From: Vishwak Thatikonda Date: Tue, 25 Aug 2026 23:43:39 -0700 Subject: [PATCH 17/43] Validate data rank in Unbatch and index rank in UnbatchGrad --- tensorflow/core/kernels/batch_kernels.cc | 14 +++++++++++ tensorflow/python/ops/batch_ops_test.py | 30 ++++++++++++++++++++++++ 2 files changed, 44 insertions(+) diff --git a/tensorflow/core/kernels/batch_kernels.cc b/tensorflow/core/kernels/batch_kernels.cc index f900c764139d77..e869c0bc55fa97 100644 --- a/tensorflow/core/kernels/batch_kernels.cc +++ b/tensorflow/core/kernels/batch_kernels.cc @@ -836,6 +836,12 @@ class UnbatchResource : public ResourceBase { "shape [batch_size, 3]; Got: ", batch_index_t.shape().DebugString(), ".")); } + if (data_t.dims() == 0) { + return absl::InvalidArgumentError( + absl::StrCat("Wrong shape for data tensor. Expected at least a " + "vector; Got: ", + data_t.shape().DebugString(), ".")); + } if (batch_index_t.shape().dim_size(0) > data_t.shape().dim_size(0)) { return absl::InvalidArgumentError(absl::StrCat( "Wrong shape for index tensor. Expected 0th dimension size to be no " @@ -1111,6 +1117,14 @@ class UnbatchGradResource : public ResourceBase { "batch_index is empty while the tensor isn't."); } std::unordered_set missing_tensors; + // The rank must be validated before any dim_size access, which has + // undefined behavior for out-of-range dimension indices. + if (!TensorShapeUtils::IsMatrix(batch_index_t.shape())) { + return absl::InvalidArgumentError( + absl::StrCat("Wrong shape for index tensor. Expected a matrix of " + "shape [batch_size, 3]; Got: ", + batch_index_t.shape().DebugString(), ".")); + } if (batch_index_t.NumElements() != batch_index_t.dim_size(0) * 3) { return absl::InvalidArgumentError(absl::StrCat( "batch_index should contain ", batch_index_t.dim_size(0) * 3, diff --git a/tensorflow/python/ops/batch_ops_test.py b/tensorflow/python/ops/batch_ops_test.py index 4c09c06d13b797..9d205e3bb4aaad 100644 --- a/tensorflow/python/ops/batch_ops_test.py +++ b/tensorflow/python/ops/batch_ops_test.py @@ -276,6 +276,36 @@ def testUnbatchInvalidIndexRank(self): container="", shared_name="") + def testUnbatchInvalidDataRank(self): + # A scalar data tensor used to reach a dimension access with undefined + # behavior and crash the process on a downstream consistency check. + if context.executing_eagerly(): + with self.assertRaisesRegex(errors.InvalidArgumentError, + "Expected at least a vector"): + batch_ops.unbatch( + batched_tensor=constant_op.constant(5.0), + batch_index=constant_op.constant([[0, 1, 0]], dtype=dtypes.int64), + id=constant_op.constant(0, dtype=dtypes.int64), + timeout_micros=0, + container="", + shared_name="") + + def testUnbatchGradInvalidIndexRank(self): + # A batch_index that is not a rank-2 matrix was rejected only by way of + # an out-of-range dimension access whose result happened to fail a later + # size comparison. Validate the rank explicitly instead. + if context.executing_eagerly(): + with self.assertRaisesRegex( + errors.InvalidArgumentError, + r"Expected a matrix of shape \[batch_size, 3\]"): + gen_batch_ops.unbatch_grad( + original_input=constant_op.constant([1.0, 2.0]), + batch_index=constant_op.constant(7, dtype=dtypes.int64), + grad=constant_op.constant([1.0, 2.0]), + id=constant_op.constant(0, dtype=dtypes.int64), + container="", + shared_name="") + def testBatchDecoratedWithCapturedInput(self): """Tests that the batch_function decorator works.""" if context.executing_eagerly(): From 116b81d7afcd1c3259083650d9da7a319b060b47 Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Tue, 5 May 2026 21:26:27 +0000 Subject: [PATCH 18/43] Fix integer overflow and type mapping in Unicode operations --- tensorflow/core/kernels/unicode_ops.cc | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/tensorflow/core/kernels/unicode_ops.cc b/tensorflow/core/kernels/unicode_ops.cc index b75e2f41e56230..657e2b4c2505e4 100644 --- a/tensorflow/core/kernels/unicode_ops.cc +++ b/tensorflow/core/kernels/unicode_ops.cc @@ -454,7 +454,7 @@ class UnicodeDecodeBaseOp : public OpKernel { Tensor* output_char_values; OP_REQUIRES_OK( ctx, ctx->allocate_output( - "char_values", {static_cast(char_values.size())}, + "char_values", {static_cast(char_values.size())}, &output_char_values)); auto out_char_values = output_char_values->vec(); if (generate_offsets_) { @@ -469,9 +469,9 @@ class UnicodeDecodeBaseOp : public OpKernel { Tensor* output_offset_values; OP_REQUIRES_OK(ctx, ctx->allocate_output( "char_to_byte_starts", - {static_cast(offset_values.size())}, + {static_cast(offset_values.size())}, &output_offset_values)); - auto out_offset_values = output_offset_values->vec(); + auto out_offset_values = output_offset_values->vec(); // Load output tensors from intermediate value arrays. for (size_t i = 0; i < char_values.size(); ++i) { From 30b8178a9ec538c221cd2526b3ed4cfec713539b Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Fri, 8 May 2026 00:23:56 +0530 Subject: [PATCH 19/43] Add missing Tsplits=int32 tests for UnicodeDecodeWithOffsets --- .../strings_ops/unicode_decode_op_test.py | 14 ++++++++++++++ 1 file changed, 14 insertions(+) diff --git a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py index eaad1f72c1fa1d..c89538a91d2808 100644 --- a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py +++ b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py @@ -463,6 +463,20 @@ def testDecodeGenOp(self, self.assertAllEqual(expected_char_to_byte_starts, result.char_to_byte_starts) + def testDecodeWithOffsetsInt32Splits(self): + """Verifies that Tsplits=dtypes.int32 processes correctly without type mismatch.""" + input_data = constant_op.constant([b"hello", b"world"], dtype=dtypes.string) + + result = gen_string_ops.unicode_decode_with_offsets( + input=input_data, + input_encoding="UTF-8", + Tsplits=dtypes.int32 + ) + + self.assertEqual(result.row_splits.dtype, dtypes.int32) + self.assertEqual(result.char_to_byte_starts.dtype, dtypes.int64) + self.assertAllEqual(result.row_splits, [0, 5, 10]) + @test_util.run_all_in_graph_and_eager_modes class UnicodeSplitTest(test_util.TensorFlowTestCase, From 5a6c469a865447ed7a52734262a8c8ad216f85a2 Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Sat, 9 May 2026 23:37:17 +0530 Subject: [PATCH 20/43] Trigger CI From 815924d22d19a2b59fb5e35a7c164228c8b3a4e6 Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Mon, 29 Jun 2026 01:23:00 +0530 Subject: [PATCH 21/43] Add value assertions for char_values and char_to_byte_starts --- .../python/kernel_tests/strings_ops/unicode_decode_op_test.py | 2 ++ 1 file changed, 2 insertions(+) diff --git a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py index c89538a91d2808..76e7db8af2db04 100644 --- a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py +++ b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py @@ -476,6 +476,8 @@ def testDecodeWithOffsetsInt32Splits(self): self.assertEqual(result.row_splits.dtype, dtypes.int32) self.assertEqual(result.char_to_byte_starts.dtype, dtypes.int64) self.assertAllEqual(result.row_splits, [0, 5, 10]) + self.assertAllEqual(result.char_values, [104, 101, 108, 108, 111, 119, 111, 114, 108, 100]) + self.assertAllEqual(result.char_to_byte_starts, [0, 1, 2, 3, 4, 0, 1, 2, 3, 4]) @test_util.run_all_in_graph_and_eager_modes From aeb4a607f78ef698980d89eecf8a43ce934c7c98 Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Tue, 30 Jun 2026 10:25:49 +0530 Subject: [PATCH 22/43] Fix pylint line-too-long errors --- .../kernel_tests/strings_ops/unicode_decode_op_test.py | 10 ++++++++-- 1 file changed, 8 insertions(+), 2 deletions(-) diff --git a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py index 76e7db8af2db04..dae7aeb060c22b 100644 --- a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py +++ b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py @@ -476,8 +476,14 @@ def testDecodeWithOffsetsInt32Splits(self): self.assertEqual(result.row_splits.dtype, dtypes.int32) self.assertEqual(result.char_to_byte_starts.dtype, dtypes.int64) self.assertAllEqual(result.row_splits, [0, 5, 10]) - self.assertAllEqual(result.char_values, [104, 101, 108, 108, 111, 119, 111, 114, 108, 100]) - self.assertAllEqual(result.char_to_byte_starts, [0, 1, 2, 3, 4, 0, 1, 2, 3, 4]) + self.assertAllEqual( + result.char_values, + [104, 101, 108, 108, 111, 119, 111, 114, 108, 100] + ) + self.assertAllEqual( + result.char_to_byte_starts, + [0, 1, 2, 3, 4, 0, 1, 2, 3, 4] + ) @test_util.run_all_in_graph_and_eager_modes From 1edb19c67753f876144b884991d9c14e7a577aa4 Mon Sep 17 00:00:00 2001 From: Abhijeet Awasthi <149562965+opabhijeet@users.noreply.github.com> Date: Wed, 1 Jul 2026 23:52:25 +0530 Subject: [PATCH 23/43] upgrade offsets to int64 and fix python formatting --- tensorflow/core/kernels/unicode_ops.cc | 6 +++--- .../kernel_tests/strings_ops/unicode_decode_op_test.py | 10 +++++----- 2 files changed, 8 insertions(+), 8 deletions(-) diff --git a/tensorflow/core/kernels/unicode_ops.cc b/tensorflow/core/kernels/unicode_ops.cc index 657e2b4c2505e4..fe0d33589a842b 100644 --- a/tensorflow/core/kernels/unicode_ops.cc +++ b/tensorflow/core/kernels/unicode_ops.cc @@ -372,7 +372,7 @@ class UnicodeDecodeBaseOp : public OpKernel { } void Decode(OpKernelContext* ctx, std::vector* char_values, - std::vector* offset_values, int* current_offset, + std::vector* offset_values, int64_t* current_offset, SPLITS_TYPE* next_row_split, UChar32 char_value, int char_length, bool found_any_format_error) { if (error_options_.error_on_malformatting && found_any_format_error) { @@ -416,7 +416,7 @@ class UnicodeDecodeBaseOp : public OpKernel { input_encoding_)); std::vector char_values; - std::vector offset_values; + std::vector offset_values; Tensor* output_row_splits; OP_REQUIRES_OK(ctx, ctx->allocate_output("row_splits", @@ -433,7 +433,7 @@ class UnicodeDecodeBaseOp : public OpKernel { // the fields needed to construct a RaggedTensor. out_row_splits(row_split_index) = next_row_split; row_split_index++; - int current_offset = 0; + int64_t current_offset = 0; IterateUnicodeString( input, input_encoder.converter_, std::bind(&UnicodeDecodeBaseOp::Decode, this, ctx, &char_values, diff --git a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py index dae7aeb060c22b..f29a8b492d8884 100644 --- a/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py +++ b/tensorflow/python/kernel_tests/strings_ops/unicode_decode_op_test.py @@ -466,13 +466,13 @@ def testDecodeGenOp(self, def testDecodeWithOffsetsInt32Splits(self): """Verifies that Tsplits=dtypes.int32 processes correctly without type mismatch.""" input_data = constant_op.constant([b"hello", b"world"], dtype=dtypes.string) - + result = gen_string_ops.unicode_decode_with_offsets( - input=input_data, - input_encoding="UTF-8", - Tsplits=dtypes.int32 + input=input_data, + input_encoding="UTF-8", + Tsplits=dtypes.int32 ) - + self.assertEqual(result.row_splits.dtype, dtypes.int32) self.assertEqual(result.char_to_byte_starts.dtype, dtypes.int64) self.assertAllEqual(result.row_splits, [0, 5, 10]) From b7f54acd78f691b657871e36a171bd4754b6c458 Mon Sep 17 00:00:00 2001 From: Volodymyr Kysenko Date: Tue, 1 Sep 2026 10:55:26 -0700 Subject: [PATCH 24/43] Fix tensorflow/lite/delegates/ynnpack/attention_bench build in open-source. PiperOrigin-RevId: 974589473 --- tensorflow/lite/delegates/ynnpack/attention_bench.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/tensorflow/lite/delegates/ynnpack/attention_bench.cc b/tensorflow/lite/delegates/ynnpack/attention_bench.cc index 050d3ea295832f..3d6c124d27cb31 100644 --- a/tensorflow/lite/delegates/ynnpack/attention_bench.cc +++ b/tensorflow/lite/delegates/ynnpack/attention_bench.cc @@ -147,7 +147,7 @@ void FullSequenceAttentionDecode(benchmark::State& state) { AttentionImpl::kFullSequence); } -void AttentionArguments(benchmark::Benchmark* b) { +void AttentionArguments(benchmark::internal::Benchmark* b) { b->ArgNames({"seq", "head", "heads", "threads", "seq_active"}); b->UseRealTime(); b->MeasureProcessCPUTime(); From cf5232b8a7c5c8a3aab5acb33bad29c50a5ea0e7 Mon Sep 17 00:00:00 2001 From: Adrian Kuegel Date: Tue, 1 Sep 2026 11:31:05 -0700 Subject: [PATCH 25/43] Reverts 7b34ce2de44d6b34e35ae10d2f5dbd7218e173a9 PiperOrigin-RevId: 974611615 --- third_party/xla/xla/debug_options_flags.cc | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/third_party/xla/xla/debug_options_flags.cc b/third_party/xla/xla/debug_options_flags.cc index 540880d59af90f..3c0289ec0c7ea0 100644 --- a/third_party/xla/xla/debug_options_flags.cc +++ b/third_party/xla/xla/debug_options_flags.cc @@ -266,7 +266,7 @@ DebugOptions DefaultDebugOptionsIgnoringFlags() { opts.set_xla_cpu_use_acl(true); #endif opts.set_xla_cpu_use_xnnpack(true); - opts.set_xla_cpu_use_new_xtile_lowering(true); + opts.set_xla_cpu_use_new_xtile_lowering(false); opts.set_xla_cpu_experimental_xnn_graph_fusion_mode( DebugOptions::XNN_GRAPH_FUSION_MODE_DISABLED); opts.add_xla_cpu_experimental_ynn_fusion_type( From 4c493e0eedb72d5d29eb4620309af1eaee770510 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Tue, 1 Sep 2026 11:50:04 -0700 Subject: [PATCH 26/43] Fix memory leak in GradientTape::ComputeGradient on backward function error. PiperOrigin-RevId: 974622549 --- tensorflow/c/eager/BUILD | 5 + tensorflow/c/eager/gradients.cc | 16 +- tensorflow/c/eager/gradients_test.cc | 200 ++++++++++++++++++++++- tensorflow/c/eager/parallel_device/BUILD | 4 - tensorflow/c/eager/tape.h | 13 +- 5 files changed, 220 insertions(+), 18 deletions(-) diff --git a/tensorflow/c/eager/BUILD b/tensorflow/c/eager/BUILD index 086afda4429adf..a4a1657545eb70 100644 --- a/tensorflow/c/eager/BUILD +++ b/tensorflow/c/eager/BUILD @@ -273,11 +273,13 @@ tf_cuda_cc_test( tags = tf_cuda_tests_tags() + ["nomac"], deps = [ ":abstract_context", + ":abstract_operation", ":abstract_tensor_handle", ":c_api_experimental", ":c_api_test_util", ":c_api_unified_internal", ":gradients_internal", + ":tape", ":unified_api_testutil", "//tensorflow/c:c_api", "//tensorflow/c:c_test_util", @@ -294,6 +296,8 @@ tf_cuda_cc_test( "//tensorflow/core:test_main", "//tensorflow/core/lib/llvm_rtti", "@com_google_absl//absl/container:flat_hash_set", + "@com_google_absl//absl/log:check", + "@com_google_absl//absl/status", "@com_google_absl//absl/strings", "@com_google_absl//absl/types:span", ], @@ -1085,6 +1089,7 @@ cc_library( "//tensorflow/core/config:flags", "//tensorflow/core/platform:errors", "//tensorflow/core/platform:types", + "@com_google_absl//absl/container:flat_hash_set", "@com_google_absl//absl/log", "@com_google_absl//absl/log:check", "@com_google_absl//absl/log:vlog_is_on", diff --git a/tensorflow/c/eager/gradients.cc b/tensorflow/c/eager/gradients.cc index 488f49a07647c4..62719bc9b60273 100644 --- a/tensorflow/c/eager/gradients.cc +++ b/tensorflow/c/eager/gradients.cc @@ -214,14 +214,22 @@ absl::Status TapeVSpace::CallBackwardFunction( const std::vector& unneeded_gradients, absl::Span output_gradients, absl::Span result) const { + absl::Status s; if (gradient_function == nullptr) { - return absl::InvalidArgumentError(absl::StrCat( + s = absl::InvalidArgumentError(absl::StrCat( "Provided null gradient_function for '", op_type, "'.\n", "If the intent is to treat this op as non-differentiable consider ", "using RegisterNotDifferentiable or ", "NotDifferentiableGradientFunction.")); + } else { + s = gradient_function->Compute(ctx_, output_gradients, result); + } + for (AbstractTensorHandle* grad : output_gradients) { + if (grad != nullptr) { + grad->Unref(); + } } - return gradient_function->Compute(ctx_, output_gradients, result); + return s; } absl::Status TapeVSpace::BuildOnesLike(const TapeTensor& t, @@ -251,7 +259,9 @@ TapeTensor TapeVSpace::TapeTensorFromGradient(AbstractTensorHandle* g) const { return TapeTensor(g); } -void TapeVSpace::MarkAsResult(AbstractTensorHandle* gradient) const {} +void TapeVSpace::MarkAsResult(AbstractTensorHandle* gradient) const { + if (gradient) gradient->Ref(); +} void TapeVSpace::DeleteGradient(AbstractTensorHandle* gradient) const { gradient->Unref(); diff --git a/tensorflow/c/eager/gradients_test.cc b/tensorflow/c/eager/gradients_test.cc index 45c4005f8204fb..c48ef3f26635de 100644 --- a/tensorflow/c/eager/gradients_test.cc +++ b/tensorflow/c/eager/gradients_test.cc @@ -14,26 +14,29 @@ limitations under the License. ==============================================================================*/ #include "tensorflow/c/eager/gradients.h" +#include #include +#include +#include -#include "absl/container/flat_hash_set.h" +#include "absl/log/check.h" +#include "absl/status/status.h" #include "absl/types/span.h" #include "tensorflow/c/eager/abstract_context.h" +#include "tensorflow/c/eager/abstract_operation.h" #include "tensorflow/c/eager/abstract_tensor_handle.h" -#include "tensorflow/c/eager/c_api_experimental.h" -#include "tensorflow/c/eager/c_api_test_util.h" #include "tensorflow/c/eager/c_api_unified_experimental.h" #include "tensorflow/c/eager/c_api_unified_experimental_internal.h" #include "tensorflow/c/eager/gradients_internal.h" +#include "tensorflow/c/eager/tape.h" #include "tensorflow/c/eager/unified_api_testutil.h" -#include "tensorflow/c/experimental/gradients/array_grad.h" -#include "tensorflow/c/experimental/gradients/math_grad.h" #include "tensorflow/c/experimental/gradients/not_differentiable.h" -#include "tensorflow/c/experimental/gradients/tape/tape_context.h" -#include "tensorflow/c/experimental/ops/array_ops.h" #include "tensorflow/c/experimental/ops/math_ops.h" +#include "tensorflow/c/tf_datatype.h" +#include "tensorflow/c/tf_status.h" #include "tensorflow/c/tf_status_helper.h" -#include "tensorflow/c/tf_tensor.h" +#include "xla/tsl/platform/errors.h" +#include "tensorflow/core/lib/gtl/array_slice.h" #include "tensorflow/core/lib/llvm_rtti/llvm_rtti.h" #include "tensorflow/core/platform/errors.h" #include "tensorflow/core/platform/test.h" @@ -123,10 +126,11 @@ absl::Status RecordOperationWithNullGradientFunctionModel( AbstractTensorHandle* neg_output; TF_RETURN_IF_ERROR(ops::Neg(ctx, inputs[0], &neg_output, "Neg")); tape.RecordOperation(inputs, {neg_output}, nullptr, "Neg"); + inputs[0]->Ref(); return tape.ComputeGradient(ctx, /*targets=*/{neg_output}, /*sources=*/inputs, - /*output_gradients=*/{}, outputs); + /*output_gradients=*/{inputs[0]}, outputs); } TEST_P(CppGradients, TestRecordOperationWithNullGradientFunctionRaises) { @@ -161,6 +165,139 @@ TEST_P(CppGradients, TestRecordOperationWithNullGradientFunctionRaises) { "or NotDifferentiableGradientFunction.", s.message()); ASSERT_EQ(nullptr, outputs[0]); + EXPECT_TRUE(x.get()->RefCountIsOne()); +} + +class DummyGradientFunction : public GradientFunction { + public: + absl::Status Compute(AbstractContext* ctx, + absl::Span grad_outputs, + absl::Span grad_inputs) override { + if (!grad_inputs.empty() && !grad_outputs.empty()) { + grad_inputs[0] = grad_outputs[0]; + if (grad_inputs[0]) { + grad_inputs[0]->Ref(); + } + } + return absl::OkStatus(); + } +}; + +TEST_P(CppGradients, TestMarkAsResult) { + std::unique_ptr status( + TF_NewStatus(), TF_DeleteStatus); + AbstractContextPtr ctx; + { + AbstractContext* ctx_raw = nullptr; + absl::Status s = + BuildImmediateExecutionContext(std::get<1>(GetParam()), &ctx_raw); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + ctx.reset(ctx_raw); + } + + AbstractTensorHandlePtr x; + { + AbstractTensorHandle* x_raw = nullptr; + absl::Status s = + TestScalarTensorHandle(ctx.get(), 2.0f, &x_raw); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + x.reset(x_raw); + } + + std::vector temp_outputs(1); + AbstractOperationPtr op(ctx.get()->CreateOperation()); + ForwardOperation forward_op; + absl::Status s = + Reset(op.get(), "Identity", /*raw_device_name=*/nullptr, &forward_op); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + s = AddInput(op.get(), x.get(), &forward_op); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + int num_retvals = 1; + s = op->Execute(absl::MakeSpan(temp_outputs), &num_retvals); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + + Tape tape(/*persistent=*/false); + tape.Watch(x.get()); + tape.RecordOperation({x.get()}, temp_outputs, new DummyGradientFunction, + "Identity"); + + std::vector outputs(1); + x.get()->Ref(); // Pass ownership of this gradient to ComputeGradient + s = tape.ComputeGradient(ctx.get(), + /*targets=*/temp_outputs, + /*sources=*/{temp_outputs[0]}, + /*output_gradients=*/{x.get()}, + absl::MakeSpan(outputs)); + ASSERT_EQ(errors::OK, s.code()); + ASSERT_EQ(x.get(), outputs[0]); + outputs[0]->Unref(); + EXPECT_TRUE(x.get()->RefCountIsOne()); +} + +struct DummyTensor { + int64_t id; + int64_t GetID() const { return id; } + tensorflow::DataType GetDType() const { return tensorflow::DT_FLOAT; } + int* ZerosLike() const { return nullptr; } +}; + +struct DummyBackwardFunction {}; + +class MockVSpace + : public eager::VSpace { + public: + mutable int delete_gradient_called_ = 0; + + int64_t NumElements(int* tensor) const override { return 1; } + int* AggregateGradients( + gtl::ArraySlice gradient_tensors) const override { + return gradient_tensors[0]; + } + absl::Status CallBackwardFunction( + const std::string& op_type, DummyBackwardFunction* backward_function, + const std::vector& unneeded_gradients, + gtl::ArraySlice output_gradients, + absl::Span result) const override { + for (int* g : output_gradients) { + if (g) DeleteGradient(g); + } + return absl::InternalError("Intentional failure"); + } + absl::Status BuildOnesLike(const DummyTensor& t, + int** result) const override { + *result = new int(1); + return absl::OkStatus(); + } + int64_t TensorId(int* tensor) const override { return 0; } + DummyTensor TapeTensorFromGradient(int* gradient) const override { + return DummyTensor{0}; + } + void MarkAsResult(int* gradient) const override {} + void DeleteGradient(int* gradient) const override { + delete_gradient_called_++; + delete gradient; + } +}; + +TEST(GradientTapeTest, MemoryLeakOnFailure) { + eager::GradientTape tape( + /*persistent=*/false); + tape.Watch(1); + + DummyBackwardFunction* bw = new DummyBackwardFunction(); + tape.RecordOperation( + "TestOp", {DummyTensor{2}}, {1}, {tensorflow::DT_FLOAT}, + [bw]() { return bw; }, [](DummyBackwardFunction* bw) { delete bw; }); + + MockVSpace vspace; + std::vector results(1); + absl::Status s = + tape.ComputeGradient(vspace, {2}, {1}, {}, {}, absl::MakeSpan(results), + /*build_default_zeros_grads=*/false); + + ASSERT_EQ(error::INTERNAL, s.code()); + EXPECT_EQ(1, vspace.delete_gradient_called_) + << "Expected gradient to be deleted (memory leak if 0)"; } TEST_P(CppGradients, TestExecuteWithLargerOutputsVectorDoesNotCrash) { @@ -235,6 +372,51 @@ INSTANTIATE_TEST_SUITE_P( /*tfrt*/ ::testing::Values(false), /*executing_eagerly*/ ::testing::Values(true, false))); #endif +TEST(GradientTapeTest, TapeVSpaceLeakOnBackwardFunctionError) { + std::unique_ptr status( + TF_NewStatus(), TF_DeleteStatus); + AbstractContextPtr ctx; + { + AbstractContext* ctx_raw = nullptr; + absl::Status s = BuildImmediateExecutionContext(false, &ctx_raw); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + ctx.reset(ctx_raw); + } + + AbstractTensorHandlePtr x; + { + AbstractTensorHandle* x_raw = nullptr; + absl::Status s = + TestScalarTensorHandle(ctx.get(), 2.0f, &x_raw); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + x.reset(x_raw); + } + + Tape tape(/*persistent=*/false); + tape.Watch(x.get()); + AbstractTensorHandle* neg_output; + absl::Status s = ops::Neg(ctx.get(), x.get(), &neg_output, "Neg"); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + AbstractTensorHandle* neg_output2; + s = ops::Neg(ctx.get(), x.get(), &neg_output2, "Neg2"); + ASSERT_EQ(errors::OK, s.code()) << s.message(); + + tape.RecordOperation({x.get()}, {neg_output, neg_output2}, nullptr, "Neg"); + x.get()->Ref(); + + std::vector outputs; + s = tape.ComputeGradient(ctx.get(), + /*targets=*/{neg_output}, + /*sources=*/{}, + /*output_gradients=*/{x.get()}, + absl::MakeSpan(outputs)); + ASSERT_EQ(error::INVALID_ARGUMENT, s.code()); + neg_output->Unref(); + neg_output2->Unref(); + + EXPECT_TRUE(x.get()->RefCountIsOne()); +} + } // namespace } // namespace internal } // namespace gradients diff --git a/tensorflow/c/eager/parallel_device/BUILD b/tensorflow/c/eager/parallel_device/BUILD index 56d7612278ec3c..559abe8c7032ff 100644 --- a/tensorflow/c/eager/parallel_device/BUILD +++ b/tensorflow/c/eager/parallel_device/BUILD @@ -79,10 +79,8 @@ cc_library( "//tensorflow/c/eager:c_api", "//tensorflow/c/eager:c_api_experimental", "//tensorflow/c/eager:tfe_tensorhandle_internal", - "//tensorflow/core/platform:status", "@com_google_absl//absl/status", "@com_google_absl//absl/strings", - "@com_google_absl//absl/types:optional", "@com_google_absl//absl/types:variant", ], ) @@ -105,7 +103,6 @@ cc_library( "//tensorflow/core:framework", "//tensorflow/core:lib", "@com_google_absl//absl/base:core_headers", - "@com_google_absl//absl/memory", "@com_google_absl//absl/status", "@com_google_absl//absl/strings", "@com_google_absl//absl/synchronization", @@ -131,7 +128,6 @@ tf_cc_test( "//tensorflow/core:test", "//tensorflow/core:test_main", "//tensorflow/core/common_runtime/eager:context", - "//tensorflow/core/platform:status", "@com_google_absl//absl/status", "@com_google_googletest//:gtest_main", ], diff --git a/tensorflow/c/eager/tape.h b/tensorflow/c/eager/tape.h index 9fc441d32d3f28..ad7c455a81309a 100644 --- a/tensorflow/c/eager/tape.h +++ b/tensorflow/c/eager/tape.h @@ -25,8 +25,10 @@ limitations under the License. #include #include #include +#include #include +#include "absl/container/flat_hash_set.h" #include "absl/log/check.h" #include "absl/log/log.h" #include "absl/log/vlog_is_on.h" @@ -812,6 +814,13 @@ GradientTape::ComputeGradient( trace.backward_function_deleter(trace.backward_function); } if (!s.ok()) { + for (const auto& pair : gradients) { + for (Gradient* g : pair.second) { + if (g != nullptr) { + vspace.DeleteGradient(g); + } + } + } return s; } } else { @@ -892,7 +901,7 @@ GradientTape::ComputeGradient( source_tensor_ids.size(), " found ", result.size(), " in call to Tape::ComputeGradient."); } - std::unordered_set used_gradient_ids(source_tensor_ids.size()); + absl::flat_hash_set used_gradient_ids(source_tensor_ids.size()); for (int i = 0; i < source_tensor_ids.size(); i++) { int64_t tensor_id = source_tensor_ids[i]; auto grad_it = gradients.find(tensor_id); @@ -967,7 +976,7 @@ ForwardAccumulator::ForwardpropFromTape( gtl::MakeCleanup([&call_state] { call_state.backward_tape = nullptr; }); std::vector forwardprop_aids; std::vector sources; - std::unordered_set sources_set; + absl::flat_hash_set sources_set; sources.reserve(output_tensors.size()); for (const TapeTensor& output_tensor : output_tensors) { // Ownership of `aid` transferred to CallBackwardFunction below. From d31f980ba594e563d72d3da28f24d11b50e3b065 Mon Sep 17 00:00:00 2001 From: Blake Hechtman Date: Tue, 1 Sep 2026 12:06:00 -0700 Subject: [PATCH 27/43] Add support for feature group convolutions (`feature_group_count > 1`) and batch group convolutions (`batch_group_count > 1`) in Mosaic TPU and JAX Pallas. This is Part 3 of a 3-part changelist chain: - Part 1 (cl/964822765): MLIR `TPU_ConvOp` dialect definition, verification, Python bindings, and basic N-D convolution lowering. - Part 2 (cl/969761722): Dilations (input/LHS and kernel/RHS), negative padding (cropping), and strided slicing emulation. - Part 3 (this CL): Grouped convolutions (`feature_group_count > 1`, `batch_group_count > 1`). PiperOrigin-RevId: 974631914 --- .../xla/xla/mosaic/dialect/tpu/tpu_ops.cc | 48 +++++-- .../xla/xla/mosaic/dialect/tpu/tpu_ops.td | 4 + .../dialect/tpu/tpu_ops_verification_test.cc | 118 +++++++++++++++++- 3 files changed, 160 insertions(+), 10 deletions(-) diff --git a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.cc b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.cc index 57afd88cfad38c..9c4a1ad856af1a 100644 --- a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.cc +++ b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.cc @@ -1432,21 +1432,47 @@ LogicalResult ConvOp::verify() { return emitOpError("Expected window attributes size to match spatial dims"); } + const int64_t feature_group_count = getFeatureGroupCount(); + const int64_t batch_group_count = getBatchGroupCount(); + if (feature_group_count <= 0) { + return emitOpError("Expected feature_group_count to be positive"); + } + if (batch_group_count <= 0) { + return emitOpError("Expected batch_group_count to be positive"); + } + if (feature_group_count > 1 && batch_group_count > 1) { + return emitOpError( + "At most one of batch_group_count and feature_group_count may be > 1"); + } + // Contracting feature dimension size match const int64_t in_feat = lhs_ty.getDimSize(dnums.getInputFeatureDimension()); const int64_t kernel_in_feat = rhs_ty.getDimSize(dnums.getKernelInputFeatureDimension()); - if (in_feat != kernel_in_feat) { + if (in_feat % feature_group_count != 0) { + return emitOpError( + absl::StrFormat("LHS feature dimension size (%d) must be divisible by " + "feature_group_count (%d)", + in_feat, feature_group_count)); + } + if (in_feat / feature_group_count != kernel_in_feat) { return emitOpError(absl::StrFormat( - "LHS feature dimension size (%d) must match kernel input feature " - "dimension size (%d)", - in_feat, kernel_in_feat)); + "LHS feature dimension size divided by feature_group_count " + "(%d / %d = %d) must match kernel input feature dimension size (%d)", + in_feat, feature_group_count, in_feat / feature_group_count, + kernel_in_feat)); } // Output feature dimension size match const int64_t out_feat = acc_ty.getDimSize(dnums.getOutputFeatureDimension()); const int64_t kernel_out_feat = rhs_ty.getDimSize(dnums.getKernelOutputFeatureDimension()); + if (kernel_out_feat % (feature_group_count * batch_group_count) != 0) { + return emitOpError(absl::StrFormat( + "Kernel output feature dimension size (%d) must be divisible by " + "feature_group_count * batch_group_count (%d)", + kernel_out_feat, feature_group_count * batch_group_count)); + } if (out_feat != kernel_out_feat) { return emitOpError(absl::StrFormat( "ACC output feature dimension size (%d) must match kernel output " @@ -1457,11 +1483,17 @@ LogicalResult ConvOp::verify() { // Batch dimension size match const int64_t in_batch = lhs_ty.getDimSize(dnums.getInputBatchDimension()); const int64_t out_batch = acc_ty.getDimSize(dnums.getOutputBatchDimension()); - if (in_batch != out_batch) { + if (in_batch % batch_group_count != 0) { + return emitOpError( + absl::StrFormat("LHS batch dimension size (%d) must be divisible by " + "batch_group_count (%d)", + in_batch, batch_group_count)); + } + if (in_batch / batch_group_count != out_batch) { return emitOpError(absl::StrFormat( - "LHS batch dimension size (%d) must match ACC output batch dimension " - "size (%d)", - in_batch, out_batch)); + "LHS batch dimension size divided by batch_group_count (%d / %d = %d) " + "must match ACC output batch dimension size (%d)", + in_batch, batch_group_count, in_batch / batch_group_count, out_batch)); } // Spatial dimension output size formula matching diff --git a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.td b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.td index ffbc06b052b25e..c9d49722e0dd9f 100644 --- a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.td +++ b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops.td @@ -909,6 +909,8 @@ def TPU_ConvOp : TPU_Op<"conv", [Pure]> { - `rhs_dilation`: Dilation factor for the kernel/filter (defaults to 1). - `window_reversal`: Boolean flags indicating whether to reverse the kernel window along each spatial dimension. + - `feature_group_count`: Number of feature groups for grouped convolution. + - `batch_group_count`: Number of batch groups for grouped convolution. - `precision`: Optional contraction precision mode for TPU matrix multiplication. }]; let arguments = (ins @@ -921,6 +923,8 @@ def TPU_ConvOp : TPU_Op<"conv", [Pure]> { DenseI64ArrayAttr:$lhs_dilation, DenseI64ArrayAttr:$rhs_dilation, DenseBoolArrayAttr:$window_reversal, + DefaultValuedAttr:$feature_group_count, + DefaultValuedAttr:$batch_group_count, OptionalAttr:$precision ); let results = (outs AnyVectorOfNonZeroRank:$result); diff --git a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops_verification_test.cc b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops_verification_test.cc index 68d8e7afa0e078..b2e04ea4e28374 100644 --- a/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops_verification_test.cc +++ b/third_party/xla/xla/mosaic/dialect/tpu/tpu_ops_verification_test.cc @@ -1465,6 +1465,8 @@ TEST_F(TpuOpsVerificationTest, ConvOpVerificationWorks) { /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(1), + /*batch_group_count=*/builder().getI64IntegerAttr(1), /*precision=*/nullptr); ASSERT_OK(VerifyOp(conv)); } @@ -1491,6 +1493,8 @@ TEST_F(TpuOpsVerificationTest, ConvOpRankMismatch) { /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(1), + /*batch_group_count=*/builder().getI64IntegerAttr(1), /*precision=*/nullptr); ASSERT_THAT(VerifyOp(conv), StatusIs(_, HasSubstr("Expected rhs rank to be 3"))); @@ -1518,11 +1522,16 @@ TEST_F(TpuOpsVerificationTest, ConvOpFeatureSizeMismatch) { /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(1), + /*batch_group_count=*/builder().getI64IntegerAttr(1), /*precision=*/nullptr); ASSERT_THAT( VerifyOp(conv), - StatusIs(_, HasSubstr("LHS feature dimension size (64) must match " - "kernel input feature dimension size (128)"))); + StatusIs( + _, + HasSubstr("LHS feature dimension size divided by feature_group_count " + "(64 / 1 = 64) must match kernel input feature dimension " + "size (128)"))); } TEST_F(TpuOpsVerificationTest, ConvOpSpatialOutputMismatch) { @@ -1547,6 +1556,8 @@ TEST_F(TpuOpsVerificationTest, ConvOpSpatialOutputMismatch) { /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(1), + /*batch_group_count=*/builder().getI64IntegerAttr(1), /*precision=*/nullptr); ASSERT_THAT(VerifyOp(conv), StatusIs(_, HasSubstr("Output spatial dimension 1 size mismatch: " @@ -1667,5 +1678,108 @@ TEST_F(TpuOpsVerificationTest, AnnotateOpVerification) { StatusIs(_, HasSubstr( "Hazard overrides are only valid for VMEM allocations"))); } + +TEST_F(TpuOpsVerificationTest, ConvOpInvalidGroupCounts) { + Value lhs = ConstantF32Vector({1, 8, 128}, {1.0f}); + Value rhs = ConstantF32Vector({3, 128, 128}, {1.0f}); + Value acc = ConstantF32Vector({1, 6, 128}, {1.0f}); + auto dnums = ConvDimensionNumbersAttr::get( + builder().getContext(), + /*input_batch_dimension=*/0, /*input_feature_dimension=*/2, + /*input_spatial_dimensions=*/{1}, + /*kernel_input_feature_dimension=*/1, + /*kernel_output_feature_dimension=*/2, + /*kernel_spatial_dimensions=*/{0}, + /*output_batch_dimension=*/0, /*output_feature_dimension=*/2, + /*output_spatial_dimensions=*/{1}); + + auto conv_invalid_fg = Create( + /*result=*/VectorType::get({1, 6, 128}, builder().getF32Type()), + /*lhs=*/lhs, /*rhs=*/rhs, /*acc=*/acc, + /*dimension_numbers=*/dnums, + /*window_strides=*/builder().getDenseI64ArrayAttr({1}), + /*padding=*/builder().getDenseI64ArrayAttr({0, 0}), + /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(0), + /*batch_group_count=*/builder().getI64IntegerAttr(1), + /*precision=*/nullptr); + ASSERT_THAT( + VerifyOp(conv_invalid_fg), + StatusIs(_, HasSubstr("Expected feature_group_count to be positive"))); + + auto conv_invalid_both = Create( + /*result=*/VectorType::get({1, 6, 128}, builder().getF32Type()), + /*lhs=*/lhs, /*rhs=*/rhs, /*acc=*/acc, + /*dimension_numbers=*/dnums, + /*window_strides=*/builder().getDenseI64ArrayAttr({1}), + /*padding=*/builder().getDenseI64ArrayAttr({0, 0}), + /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(2), + /*batch_group_count=*/builder().getI64IntegerAttr(2), + /*precision=*/nullptr); + ASSERT_THAT(VerifyOp(conv_invalid_both), + StatusIs(_, HasSubstr("At most one of batch_group_count and " + "feature_group_count may be > 1"))); +} + +TEST_F(TpuOpsVerificationTest, ConvOpFeatureGroupCountValid) { + Value lhs = ConstantF32Vector({1, 8, 128}, {1.0f}); + Value rhs = ConstantF32Vector({3, 64, 128}, {1.0f}); + Value acc = ConstantF32Vector({1, 6, 128}, {1.0f}); + auto dnums = ConvDimensionNumbersAttr::get( + builder().getContext(), + /*input_batch_dimension=*/0, /*input_feature_dimension=*/2, + /*input_spatial_dimensions=*/{1}, + /*kernel_input_feature_dimension=*/1, + /*kernel_output_feature_dimension=*/2, + /*kernel_spatial_dimensions=*/{0}, + /*output_batch_dimension=*/0, /*output_feature_dimension=*/2, + /*output_spatial_dimensions=*/{1}); + auto conv = Create( + /*result=*/VectorType::get({1, 6, 128}, builder().getF32Type()), + /*lhs=*/lhs, /*rhs=*/rhs, /*acc=*/acc, + /*dimension_numbers=*/dnums, + /*window_strides=*/builder().getDenseI64ArrayAttr({1}), + /*padding=*/builder().getDenseI64ArrayAttr({0, 0}), + /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(2), + /*batch_group_count=*/builder().getI64IntegerAttr(1), + /*precision=*/nullptr); + ASSERT_OK(VerifyOp(conv)); +} + +TEST_F(TpuOpsVerificationTest, ConvOpBatchGroupCountValid) { + Value lhs = ConstantF32Vector({4, 8, 128}, {1.0f}); + Value rhs = ConstantF32Vector({3, 128, 64}, {1.0f}); + Value acc = ConstantF32Vector({2, 6, 64}, {1.0f}); + auto dnums = ConvDimensionNumbersAttr::get( + builder().getContext(), + /*input_batch_dimension=*/0, /*input_feature_dimension=*/2, + /*input_spatial_dimensions=*/{1}, + /*kernel_input_feature_dimension=*/1, + /*kernel_output_feature_dimension=*/2, + /*kernel_spatial_dimensions=*/{0}, + /*output_batch_dimension=*/0, /*output_feature_dimension=*/2, + /*output_spatial_dimensions=*/{1}); + auto conv = Create( + /*result=*/VectorType::get({2, 6, 64}, builder().getF32Type()), + /*lhs=*/lhs, /*rhs=*/rhs, /*acc=*/acc, + /*dimension_numbers=*/dnums, + /*window_strides=*/builder().getDenseI64ArrayAttr({1}), + /*padding=*/builder().getDenseI64ArrayAttr({0, 0}), + /*lhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*rhs_dilation=*/builder().getDenseI64ArrayAttr({1}), + /*window_reversal=*/builder().getDenseBoolArrayAttr({false}), + /*feature_group_count=*/builder().getI64IntegerAttr(1), + /*batch_group_count=*/builder().getI64IntegerAttr(2), + /*precision=*/nullptr); + ASSERT_OK(VerifyOp(conv)); +} } // namespace } // namespace mlir::tpu From 8eee66319446737cdb639979c88a4a62903eba63 Mon Sep 17 00:00:00 2001 From: Pavel Emeliyanenko Date: Tue, 1 Sep 2026 12:07:40 -0700 Subject: [PATCH 28/43] PR #47176: [ROCM] Wire all MORI collectives through a dedicated CollectivesFacade Imported from GitHub PR https://github.com/openxla/xla/pull/47176 ## Summary This PR routes every MORI GPU collective in `MoriCommunicator` through a single, dedicated `CollectivesFacade` entry point instead of the previous ad-hoc per-op plumbing. It also introduces a thin shared header and an inert stub facade so the XLA side compiles, links, and can be committed **without** depending on the external `@roc_mori` library. The facade exposes a **non-templated, enum-based** public API: `Run*` methods take `mori::collective::DataType` and `ReduceOpKind` enums rather than `` template arguments. Host translation units (the communicator) see decl-only method declarations, while a single device translation unit (`mori_kernels.cu.cc`, compiled as HIP) defines `MORI_KERNELS_IMPL` before including the facade and thereby compiles the method definitions exactly once. This replaces the earlier `template ` + explicit-instantiation machinery. ## Motivation - Consolidate all collective launches (all-reduce, reduce-scatter, all-gather, all-to-all, barrier, send/recv, collective-permute, quiet, fence) behind one facade owned by the communicator, so staging buffers and group counters have a single, clearly-scoped owner and lifetime. - Present a clean compiled boundary: a small non-templated public API instead of template dispatch leaking through the header, no explicit-instantiation lists to maintain, and the type/op dispatch localized inside the facade. - Decouple the XLA build from the MORI source tree so this wiring can land independently, with a compile-time switch to bring in the real device facade later. ## What changed ### `mori_communicator.cc` / `mori_communicator.h` - Each communicator owns a `std::unique_ptr`, created in `MoriCommunicator::Create` (records rank identity and allocates the symmetric-heap staging buffer). The `unique_ptr` frees the staging/counters via the facade destructor before `ShmemFinalize`. - All `Launch*` paths now call `facade_->Run*` and convert the returned `hipError_t` to `absl::Status` via `se::gpu::ToStatus`. - Reduction dispatch is now enum translation, not macro/switch-key generation: `ToMoriDataType(PrimitiveType)` and `ToMoriReduceOp(ReductionKind)` map XLA enums to the facade enums (returning `Unimplemented` for unsupported dtypes), then `LaunchAllReduce` / `LaunchReduceScatter` call the non-templated `facade_->RunAllReduce(...)` / `RunReduceScatter(...)` directly. The old `MoriRedKey` key-packing and `MORI_FOR_EACH_DTYPE` x `MORI_FOR_EACH_OP` case tables are gone. - Supported dtypes include the OCP fp8 types `F8E5M2` and `F8E4M3FN` alongside F16/BF16/S8/U8/S32/U32/S64/U64/F32/F64. - Removed the old `P2P`/`P2PType` indirection; `Send`/`Recv` now go straight through `LaunchSend`/`LaunchRecv`. `Send` dropped its unused `recv_buffer` parameter and `Recv` its unused `send_buffer` parameter. - Stream handle helper switched from `AsRocmStream` (intptr) to `AsHipStream` (`hipStream_t`); `ToStream` moved into an anonymous namespace. - `MoriCommunicator::Create` now validates `num_ranks > 0`. ### `mori_kernels.h` (new) - Thin seam included by both the host communicator (decl-only) and the device TU. It selects the facade implementation: by default includes the inert stub (`mori_stub.h`); defining `XLA_GPU_USE_REAL_MORI` pulls in the real `mori/collective/collectives_facade.hpp` instead. - No longer defines any dtype/op expansion macros; the type/op dispatch now lives inside the facade. ### `mori_kernels.cu.cc` (new) - Single HIP device TU. It `#define`s `MORI_KERNELS_IMPL` and includes `mori_kernels.h`, which compiles the facade's device path (kernels + non-templated `Run*` definitions) exactly once and emits the symbols the host references. No explicit template instantiations. In the default (stub) build it is a harmless near-empty TU. ### `mori_stub.h` - Inert, header-only `mori::collective::CollectivesFacade` mirroring the real facade's non-templated enum API. Defines the `DataType` (incl. `F8E5M2`/`F8E4M3FN`) and `ReduceOpKind` enums and the `AddressVector` alias. `Create` returns a valid empty facade; every `Run*` is a no-op returning `hipSuccess`. Copybara import of the project: -- a498a2ead52c3aba6721f5edd85d48b6faa6dce3 by Pavel Emeliyanenko : wired all MORI collectives through dedicated collectives facade added symm mem back fix after rebase update collective facade usage fixing absl macros large changes in mori communicator refactoring mori communicator cosmetics fixing build switched to enum-based interface updated collective permute interface fixed clang tidy Revert "fixed clang tidy" This reverts commit f9750eafa37dd81f55c1fddb19de2b0765942217. Revert "updated collective permute interface" This reverts commit 34ca3ef080dd970e82e15bbf65358c0e8da15d8a. adapted interface -- 47a25ee3e2d04ea54c4f41b4deb8ec3b393bb6f7 by Pavel Emeliyanenko : update after rebase Merging this change closes #47176 PiperOrigin-RevId: 974632888 --- .../xla/xla/backends/gpu/collectives/BUILD | 36 ++- .../gpu/collectives/mori_communicator.cc | 286 ++++++++++++------ .../gpu/collectives/mori_communicator.h | 79 ++--- .../gpu/collectives/mori_kernels.cu.cc | 19 ++ .../backends/gpu/collectives/mori_kernels.h | 27 ++ .../xla/backends/gpu/collectives/mori_stub.h | 75 +++++ 6 files changed, 390 insertions(+), 132 deletions(-) create mode 100644 third_party/xla/xla/backends/gpu/collectives/mori_kernels.cu.cc create mode 100644 third_party/xla/xla/backends/gpu/collectives/mori_kernels.h diff --git a/third_party/xla/xla/backends/gpu/collectives/BUILD b/third_party/xla/xla/backends/gpu/collectives/BUILD index 6e6aadc1c6ff5e..b21fc137166589 100644 --- a/third_party/xla/xla/backends/gpu/collectives/BUILD +++ b/third_party/xla/xla/backends/gpu/collectives/BUILD @@ -1,4 +1,8 @@ -load("@local_config_rocm//rocm:build_defs.bzl", "if_rocm_is_configured") +load( + "@local_config_rocm//rocm:build_defs.bzl", + "if_rocm_is_configured", + "rocm_library", +) load("@local_config_sycl//sycl:build_defs.bzl", "if_sycl_is_configured") load("//xla:xla.default.bzl", "xla_cc_test") load("//xla/stream_executor:build_defs.bzl", "if_cuda_or_rocm_is_configured") @@ -1267,6 +1271,26 @@ cc_library( ], ) +rocm_library( + name = "mori_kernels", + srcs = [ + "mori_kernels.cu.cc", + ], + hdrs = [ + "mori_kernels.h", + "mori_stub.h", + ], + # copybara:uncomment compatible_with = ["//buildenv/target:non_prod"], + copts = ["-U__HIP_DISABLE_CPP_FUNCTIONS__"], # <-- only if needed + linkstatic = True, + tags = [ + "gpu", + "no-oneapi", + "rocm-only", + ], + deps = [], +) + cc_library( name = "mori_collectives", srcs = [ @@ -1276,10 +1300,13 @@ cc_library( hdrs = [ "mori_collectives.h", "mori_communicator.h", + "mori_kernels.h", "mori_stub.h", ], tags = [ "gpu", + "no-oneapi", + "rocm-only", ], visibility = ["//visibility:public"], deps = [ @@ -1299,6 +1326,7 @@ cc_library( "//xla/core/collectives:communicator", "//xla/core/collectives:rank_id", "//xla/core/collectives:reduction_kind", + "//xla/core/collectives:symmetric_memory", "//xla/pjrt/distributed:key_value_store_interface", "//xla/runtime:device_id", "//xla/runtime:process_id", @@ -1307,6 +1335,7 @@ cc_library( "//xla/stream_executor:platform_manager", "//xla/stream_executor:stream", "//xla/stream_executor:stream_executor_h", + "//xla/stream_executor/rocm:rocm_status", "//xla/tsl/platform:env", "@com_google_absl//absl/algorithm:container", "@com_google_absl//absl/base", @@ -1329,7 +1358,10 @@ cc_library( "@com_google_absl//absl/types:span", "@tsl//tsl/platform:casts", "@tsl//tsl/platform:numbers", - ], + ] + if_rocm_is_configured([ + ":mori_kernels", + "@local_config_rocm//rocm:rocm_headers", + ]), alwayslink = True, ) diff --git a/third_party/xla/xla/backends/gpu/collectives/mori_communicator.cc b/third_party/xla/xla/backends/gpu/collectives/mori_communicator.cc index d1d71dc7fd55c8..8dfb3bfa4e80e9 100644 --- a/third_party/xla/xla/backends/gpu/collectives/mori_communicator.cc +++ b/third_party/xla/xla/backends/gpu/collectives/mori_communicator.cc @@ -36,13 +36,14 @@ limitations under the License. #include "xla/backends/gpu/collectives/cancellation_token.h" #include "xla/backends/gpu/collectives/gpu_collectives.h" #include "xla/backends/gpu/collectives/mori_collectives.h" -#include "xla/backends/gpu/collectives/mori_stub.h" +#include "xla/backends/gpu/collectives/mori_kernels.h" #include "xla/core/collectives/communicator.h" #include "xla/core/collectives/rank_id.h" #include "xla/core/collectives/reduction_kind.h" #include "xla/future.h" #include "xla/primitive_util.h" #include "xla/stream_executor/device_address.h" +#include "xla/stream_executor/rocm/rocm_status.h" #include "xla/stream_executor/stream.h" #include "xla/util.h" #include "xla/xla_data.pb.h" @@ -51,31 +52,104 @@ limitations under the License. namespace shmem = ::mori::shmem; namespace xla::gpu { -static auto AsRocmStream(se::Stream* stream) { - return reinterpret_cast( +using ::mori::collective::CollectivesFacade; +namespace { + +hipStream_t AsHipStream(se::Stream* stream) { + return reinterpret_cast( stream->platform_specific_handle().stream); } -static size_t ToMoriByteCount(PrimitiveType dtype, size_t count) { +size_t ToMoriByteCount(PrimitiveType dtype, size_t count) { if (primitive_util::IsComplexType(dtype)) { count *= 2; } return count * primitive_util::BitWidth(dtype) / 8; } +absl::StatusOr<::mori::collective::DataType> ToMoriDataType( + PrimitiveType dtype) { +#define MORI_TYPE_DISPATCH(x) \ + case x: \ + return ::mori::collective::DataType::x; + switch (dtype) { + MORI_TYPE_DISPATCH(F8E5M2) + MORI_TYPE_DISPATCH(F8E4M3FN) + MORI_TYPE_DISPATCH(F16) + MORI_TYPE_DISPATCH(BF16) + MORI_TYPE_DISPATCH(S8) + MORI_TYPE_DISPATCH(U8) + MORI_TYPE_DISPATCH(S32) + MORI_TYPE_DISPATCH(U32) + MORI_TYPE_DISPATCH(S64) + MORI_TYPE_DISPATCH(U64) + MORI_TYPE_DISPATCH(F32) + MORI_TYPE_DISPATCH(F64) + default: + return absl::UnimplementedError(absl::StrFormat( + "MORI: unsupported dtype: %d", static_cast(dtype))); + } +#undef MORI_TYPE_DISPATCH +} + +// Translate an XLA ReductionKind to the facade's reduction-op enum. +absl::StatusOr<::mori::collective::ReduceOpKind> ToMoriReduceOp( + ReductionKind r) { +#define MORI_OP_DISPATCH(x) \ + case ReductionKind::x: \ + return ::mori::collective::ReduceOpKind::x; + switch (r) { + MORI_OP_DISPATCH(SUM) + MORI_OP_DISPATCH(PRODUCT) + MORI_OP_DISPATCH(MIN) + MORI_OP_DISPATCH(MAX) + default: + return absl::UnimplementedError(absl::StrFormat( + "MORI: unsupported reduction op: %d", static_cast(r))); + } +#undef MORI_OP_DISPATCH +} + +absl::StatusOr ToStream(const Communicator::Executor& executor) { + if (auto* gpu_executor = + absl::down_cast(&executor)) { + return gpu_executor->stream(); + } + return InvalidArgument("Communicator executor is not a GPU executor"); +} +} // namespace + absl::StatusOr> MoriCommunicator::Create( MoriCollectives* coll, std::shared_ptr cancel, int rank, absl::Span rank_to_pe) { auto comm = absl::WrapUnique(new MoriCommunicator(coll, cancel)); const int num_ranks = static_cast(rank_to_pe.size()); + if (num_ranks <= 0) { + return absl::InvalidArgumentError(absl::StrFormat( + "MoriCommunicator: unsupported number of ranks %d", num_ranks)); + } comm->rank_ = rank; comm->num_ranks_ = num_ranks; + + // The CollectivesFacade owns this communicator's symmetric-heap staging + // buffer and the push reduce-scatter group counters. It records the rank + // identity (rank/num_ranks) and allocates the ~2GB staging; the unique_ptr + // frees it (before ShmemFinalize) when the communicator is destroyed. + const size_t buffer_size = 2UL << 30; // 2GB + comm->facade_ = CollectivesFacade::Create(rank, num_ranks, buffer_size); + if (comm->facade_ == nullptr) { + return absl::InternalError("CollectivesFacade::Create failed"); + } VLOG(1) << "Created " << *comm << " with participants: " << num_ranks; return comm; } -MoriCommunicator::~MoriCommunicator() {} +MoriCommunicator::~MoriCommunicator() { + // facade_ (unique_ptr) releases this communicator's staging + counters via + // the CollectivesFacade dtor here, before MoriCollectives::Finalize() -> + // ShmemFinalize. +} #define CHECK_CANCELLED() \ if (cancel_->IsCancelled()) { \ @@ -97,41 +171,26 @@ absl::Status MoriCommunicator::Abort() { return absl::OkStatus(); } -absl::Status MoriCommunicator::Barrier(const Communicator::Executor& executor) { +absl::Status MoriCommunicator::Barrier(const Executor& executor) { VLOG(1) << "Barrier: " << ToString(); CHECK_CANCELLED() ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - (void)stream; - // return xla_mori::BarrierOnStream(AsRocmStream(stream)); - return absl::OkStatus(); + return se::gpu::ToStatus(facade_->RunBarrier(AsHipStream(stream))); } absl::StatusOr MoriCommunicator::NumRanks() const { - VLOG(5) << "Get the number of ranks in MORI communicator: " << ToString(); CHECK_CANCELLED() - return static_cast(num_ranks_); } absl::StatusOr MoriCommunicator::CurrentRank() { - VLOG(5) << "Get current rank in MORI communicator: " << ToString(); CHECK_CANCELLED() - return static_cast(rank_); } std::string MoriCommunicator::ToString() const { - return absl::StrFormat("MoriCommunicator(rank=%d, num_ranks=%d, my_pe=%d)", - rank_, num_ranks_, shmem::ShmemMyPe()); -} - -absl::StatusOr MoriCommunicator::ToStream( - const Executor& executor) { - if (auto* gpu_executor = - absl::down_cast(&executor)) { - return gpu_executor->stream(); - } - return InvalidArgument("Communicator executor is not a GPU executor"); + return absl::StrFormat("MoriCommunicator(rank=%d, num_ranks=%d)", rank_, + num_ranks_); } Future<> MoriCommunicator::AllReduce(se::DeviceAddressBase send_buffer, @@ -201,20 +260,20 @@ Future<> MoriCommunicator::CollectivePermute( }); } -Future<> MoriCommunicator::Send(se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, +Future<> MoriCommunicator::Send(se::DeviceAddressBase send_buffer, PrimitiveType dtype, size_t count, RankId peer, const Executor& executor) { - return P2P(P2PType::Send, dtype, recv_buffer, send_buffer, count, peer, - executor); + return Execute([send_buffer, dtype, count, peer, &executor, this]() { + return LaunchSend(send_buffer, dtype, count, peer, executor); + }); } Future<> MoriCommunicator::Recv(se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, PrimitiveType dtype, size_t count, RankId peer, const Executor& executor) { - return P2P(P2PType::Recv, dtype, recv_buffer, send_buffer, count, peer, - executor); + return Execute([recv_buffer, dtype, count, peer, &executor, this]() { + return LaunchRecv(recv_buffer, dtype, count, peer, executor); + }); } absl::Status MoriCommunicator::LaunchAllGather( @@ -222,11 +281,14 @@ absl::Status MoriCommunicator::LaunchAllGather( PrimitiveType dtype, size_t count, const Executor& executor) { CHECK_CANCELLED() ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - VLOG(3) << "LaunchAllGather: send_buffer=" << send_buffer.opaque() + + VLOG(3) << "Launch AllGather: send_buffer=" << send_buffer.opaque() << " recv_buffer=" << recv_buffer.opaque() << " count=" << count << " dtype=" << primitive_util::LowercasePrimitiveTypeName(dtype) - << " stream=" << AsRocmStream(stream); - return absl::UnimplementedError("Not implemented"); + << " stream=" << AsHipStream(stream); + return se::gpu::ToStatus(facade_->RunAllGather( + send_buffer.opaque(), recv_buffer.opaque(), ToMoriByteCount(dtype, count), + AsHipStream(stream))); } absl::Status MoriCommunicator::LaunchAllReduce( @@ -234,25 +296,20 @@ absl::Status MoriCommunicator::LaunchAllReduce( PrimitiveType dtype, size_t count, ReductionKind reduction_kind, const Executor& executor) { CHECK_CANCELLED() - ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - auto gpu_stream = AsRocmStream(stream); - (void)gpu_stream; - void* source_ptr = send_buffer.opaque(); - void* dest_ptr = recv_buffer.opaque(); - (void)source_ptr; - (void)dest_ptr; - if (primitive_util::IsComplexType(dtype)) { - count *= 2; - } VLOG(3) << absl::StreamFormat( - "Launch MORI AllReduce send_buffer=%p; recv_buffer=%p; dtype=%s; " - "count=%d; reduction_kind=%v; device_ordinal=%d", + "Launch AllReduce: send_buffer=%p; recv_buffer=%p; dtype=%s; count=%d; " + "reduction_kind=%v; stream=%p", send_buffer.opaque(), recv_buffer.opaque(), primitive_util::LowercasePrimitiveTypeName(dtype), count, reduction_kind, - stream->parent()->device_ordinal()); - return absl::UnimplementedError("Not implemented"); + stream); + + ABSL_ASSIGN_OR_RETURN(auto dt, ToMoriDataType(dtype)); + ABSL_ASSIGN_OR_RETURN(auto op, ToMoriReduceOp(reduction_kind)); + return se::gpu::ToStatus(facade_->RunAllReduce(send_buffer.opaque(), + recv_buffer.opaque(), count, + dt, op, AsHipStream(stream))); } absl::Status MoriCommunicator::LaunchReduceScatter( @@ -265,8 +322,78 @@ absl::Status MoriCommunicator::LaunchReduceScatter( VLOG(3) << "LaunchReduceScatter: send_buffer=" << send_buffer.opaque() << " recv_buffer=" << recv_buffer.opaque() << " count=" << count << " dtype=" << primitive_util::LowercasePrimitiveTypeName(dtype) - << " stream=" << AsRocmStream(stream); - return absl::UnimplementedError("Not implemented"); + << " stream=" << AsHipStream(stream); + + ABSL_ASSIGN_OR_RETURN(auto dt, ToMoriDataType(dtype)); + ABSL_ASSIGN_OR_RETURN(auto op, ToMoriReduceOp(kind)); + return se::gpu::ToStatus( + facade_->RunReduceScatter(send_buffer.opaque(), recv_buffer.opaque(), + count, dt, op, AsHipStream(stream))); +} + +absl::Status MoriCommunicator::LaunchAllToAll( + absl::InlinedVector send_buffers, + absl::InlinedVector recv_buffers, + PrimitiveType dtype, size_t count, const Executor& executor) { + CHECK_CANCELLED() + ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); + + auto format_addr = [](std::string* out, se::DeviceAddressBase buf) { + absl::StrAppendFormat(out, "%p", buf.opaque()); + }; + VLOG(3) << absl::StreamFormat( + "Launch MORI AllToAll operation; send_buffers=[%s]; recv_buffers=[%s]; " + "dtype=%s; count=%d; stream=%p", + absl::StrJoin(send_buffers, ", ", format_addr), + absl::StrJoin(recv_buffers, ", ", format_addr), + primitive_util::LowercasePrimitiveTypeName(dtype), count, + AsHipStream(stream)); + + if (send_buffers.size() != recv_buffers.size() || + send_buffers.size() != static_cast(num_ranks_)) { + return InvalidArgument( + "Number of send/recv buffers and number of ranks mismatch"); + } + + CollectivesFacade::AddressVector addrs; + addrs.reserve(num_ranks_); + for (int p = 0; p < num_ranks_; ++p) { + addrs.emplace_back(send_buffers[p].opaque(), recv_buffers[p].opaque()); + } + return se::gpu::ToStatus(facade_->RunAllToAll( + addrs, ToMoriByteCount(dtype, count), AsHipStream(stream))); +} + +absl::Status MoriCommunicator::LaunchSend(se::DeviceAddressBase send_buffer, + PrimitiveType dtype, size_t count, + RankId peer, + const Executor& executor) { + CHECK_CANCELLED() + ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); + VLOG(3) << absl::StreamFormat( + "Launch MORI Send operation; send_buffer=%p; dtype=%s; count=%d; " + "peer=%d; stream=%p", + send_buffer.opaque(), primitive_util::LowercasePrimitiveTypeName(dtype), + count, peer.value(), AsHipStream(stream)); + return se::gpu::ToStatus( + facade_->RunSend(send_buffer.opaque(), ToMoriByteCount(dtype, count), + static_cast(peer.value()), AsHipStream(stream))); +} + +absl::Status MoriCommunicator::LaunchRecv(se::DeviceAddressBase recv_buffer, + PrimitiveType dtype, size_t count, + RankId peer, + const Executor& executor) { + CHECK_CANCELLED() + ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); + VLOG(3) << absl::StreamFormat( + "Launch MORI Recv operation; recv_buffer=%p; dtype=%s; count=%d; " + "peer=%d; stream=%p", + recv_buffer.opaque(), primitive_util::LowercasePrimitiveTypeName(dtype), + count, peer.value(), AsHipStream(stream)); + return se::gpu::ToStatus( + facade_->RunRecv(recv_buffer.opaque(), ToMoriByteCount(dtype, count), + static_cast(peer.value()), AsHipStream(stream))); } absl::Status MoriCommunicator::LaunchCollectivePermute( @@ -275,13 +402,11 @@ absl::Status MoriCommunicator::LaunchCollectivePermute( absl::Span target_ranks, const Executor& executor) { CHECK_CANCELLED() ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - size_t bytes = ToMoriByteCount(dtype, count); - (void)bytes; auto rank_formatter = [](std::string* out, RankId rank) { absl::StrAppendFormat(out, "%d", rank.value()); }; VLOG(3) << absl::StreamFormat( - "[%d] Launch MORI CollectivePermute operation; send_buffer=%p; " + "[%d] Launch CollectivePermute: send_buffer=%p; " "recv_buffer=%p; dtype=%s; source_rank=%s; target_[ranks=%s]; count=%d; " "stream=%p", stream->parent()->device_ordinal(), send_buffer.opaque(), @@ -289,41 +414,15 @@ absl::Status MoriCommunicator::LaunchCollectivePermute( source_rank ? absl::StrCat(source_rank->value()) : "", absl::StrJoin(target_ranks, ", ", rank_formatter), count, stream); - return absl::UnimplementedError("Not implemented"); -} - -// Performs point-to-point communication between two ranks using MORI. -// Send: launches a single GPU kernel that copies data to the peer via P2P -// and sets a completion flag on the peer. -// Recv: launches a single-thread GPU kernel that waits for the flag. -absl::Status MoriCommunicator::P2P(P2PType p2p_type, PrimitiveType dtype, - se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, - size_t count, RankId peer, - const Executor& executor) { - const char* stype = (p2p_type == P2PType::Send ? " Send" : " Recv"); - VLOG(1) << CurrentRank().value() << stype << " to " << peer.value() - << " count " << count << " MORI communicator: " << ToString(); - CHECK_CANCELLED() - - void* source_ptr = send_buffer.opaque(); - void* dest_ptr = recv_buffer.opaque(); - - ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - auto gpu_stream = AsRocmStream(stream); - size_t bytes = ToMoriByteCount(dtype, count); - int res = 0; - (void)bytes; - (void)res; - (void)gpu_stream; - (void)source_ptr; - (void)dest_ptr; - (void)peer; - (void)stream; - (void)dtype; - (void)count; - (void)p2p_type; - return absl::UnimplementedError("Not implemented"); + std::vector dstPes; + dstPes.reserve(target_ranks.size()); + for (RankId rank : target_ranks) { + dstPes.push_back(static_cast(rank.value())); + } + const int srcPe = source_rank ? static_cast(source_rank->value()) : -1; + return se::gpu::ToStatus(facade_->RunCollectivePermute( + send_buffer.opaque(), recv_buffer.opaque(), ToMoriByteCount(dtype, count), + srcPe, dstPes, AsHipStream(stream))); } Future<> MoriCommunicator::GroupExecute( @@ -339,19 +438,16 @@ absl::Status MoriCommunicator::GroupLaunch( } absl::Status MoriCommunicator::Quiet(const Executor& executor) { - VLOG(1) << "Quiet MORI communicator: " << ToString(); + VLOG(1) << "Quiet: " << ToString(); CHECK_CANCELLED() ABSL_ASSIGN_OR_RETURN(se::Stream * stream, ToStream(executor)); - auto gpu_stream = AsRocmStream(stream); - (void)gpu_stream; - return absl::UnimplementedError("Not implemented"); + return se::gpu::ToStatus(facade_->RunQuiet(AsHipStream(stream))); } absl::Status MoriCommunicator::Fence() { - VLOG(1) << "Fence MORI communicator: " << ToString(); + VLOG(1) << "Fence: " << ToString(); CHECK_CANCELLED() - // rocm_mori_fence(); - return absl::UnimplementedError("Not implemented"); + return se::gpu::ToStatus(facade_->RunFence()); } absl::Status MoriCommunicator::PollUntilDone() const { diff --git a/third_party/xla/xla/backends/gpu/collectives/mori_communicator.h b/third_party/xla/xla/backends/gpu/collectives/mori_communicator.h index ac4b2ffb0e88b7..61b7f14a12c394 100644 --- a/third_party/xla/xla/backends/gpu/collectives/mori_communicator.h +++ b/third_party/xla/xla/backends/gpu/collectives/mori_communicator.h @@ -25,6 +25,7 @@ limitations under the License. #include "absl/functional/function_ref.h" #include "absl/status/status.h" #include "absl/status/statusor.h" +#include "absl/strings/str_format.h" #include "absl/strings/string_view.h" #include "absl/types/span.h" #include "xla/backends/gpu/collectives/cancellation_token.h" @@ -32,20 +33,44 @@ limitations under the License. #include "xla/core/collectives/communicator.h" #include "xla/core/collectives/rank_id.h" #include "xla/core/collectives/reduction_kind.h" +#include "xla/core/collectives/symmetric_memory.h" #include "xla/future.h" #include "xla/stream_executor/device_address.h" #include "xla/stream_executor/stream.h" #include "xla/xla_data.pb.h" +namespace mori::collective { +class CollectivesFacade; +} // namespace mori::collective + namespace xla::gpu { class MoriCollectives; +// Dummy symmetric memory for the MORI backend. MORI collective buffers are +// allocated directly from the symmetric shmem heap, so the local device +// address doubles as the symmetric handle and no separate registration is +// required. This simply returns the address it was created with. +class MoriSymmetricMemory : public SymmetricMemory { + public: + explicit MoriSymmetricMemory(se::DeviceAddressBase addr) : addr_(addr) {} + + se::DeviceAddressBase addr() const final { return addr_; } + + std::string ToString() const final { + return absl::StrFormat("MoriSymmetricMemory(addr=%p, size=%d)", + addr_.opaque(), addr_.size()); + } + + PackedKernelArg PackKernelArg() const final { return addr_.opaque(); } + + private: + se::DeviceAddressBase addr_; +}; + // XLA collectives communicator wrapping a MORI communicator. class MoriCommunicator : public GpuCommunicator { public: - constexpr static uint32_t kMaxTeams = 24; - friend class MoriCollectives; ~MoriCommunicator() override; @@ -63,6 +88,14 @@ class MoriCommunicator : public GpuCommunicator { absl::StatusOr NumRanks() const final; absl::StatusOr CurrentRank() final; + absl::StatusOr> CreateSymmetricMemory( + se::DeviceAddressBase addr) final { + // Dummy implementation: MORI buffers are already allocated from the + // symmetric shmem heap, so the local device address is the symmetric + // handle. Just wrap and return it unchanged. + return std::make_unique(addr); + } + absl::Status Barrier(const Executor& executor) final; Future<> GroupExecute(absl::AnyInvocable group) final; @@ -98,21 +131,9 @@ class MoriCommunicator : public GpuCommunicator { const Executor& executor) final; Future<> Send(se::DeviceAddressBase send_buffer, PrimitiveType dtype, - size_t count, RankId peer, const Executor& executor) final { - return absl::UnimplementedError("Not implemented"); - } - - Future<> Recv(se::DeviceAddressBase recv_buffer, PrimitiveType dtype, - size_t count, RankId peer, const Executor& executor) final { - return absl::UnimplementedError("Not implemented"); - } - - Future<> Send(se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, PrimitiveType dtype, size_t count, RankId peer, const Executor& executor) final; - Future<> Recv(se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, PrimitiveType dtype, + Future<> Recv(se::DeviceAddressBase recv_buffer, PrimitiveType dtype, size_t count, RankId peer, const Executor& executor) final; // Polls the communicator until any pending non-blocking operations are done @@ -149,9 +170,7 @@ class MoriCommunicator : public GpuCommunicator { absl::Status LaunchAllToAll( absl::InlinedVector send_buffers, absl::InlinedVector recv_buffers, - PrimitiveType dtype, size_t count, const Executor& executor) final { - return absl::UnimplementedError("Not implemented"); - } + PrimitiveType dtype, size_t count, const Executor& executor) final; absl::Status LaunchCollectivePermute(se::DeviceAddressBase send_buffer, se::DeviceAddressBase recv_buffer, @@ -162,15 +181,11 @@ class MoriCommunicator : public GpuCommunicator { absl::Status LaunchSend(se::DeviceAddressBase send_buffer, PrimitiveType dtype, size_t count, RankId peer, - const Executor& executor) final { - return absl::UnimplementedError("Not implemented"); - } + const Executor& executor) final; absl::Status LaunchRecv(se::DeviceAddressBase recv_buffer, PrimitiveType dtype, size_t count, RankId peer, - const Executor& executor) final { - return absl::UnimplementedError("Not implemented"); - } + const Executor& executor) final; absl::Status Quiet(const Executor& executor) final; @@ -186,22 +201,16 @@ class MoriCommunicator : public GpuCommunicator { std::shared_ptr cancel) : collectives_(coll), cancel_(std::move(cancel)) {} - enum class P2PType : int32_t { Send, Recv }; - - absl::Status P2P(P2PType p2p_type, PrimitiveType type, - se::DeviceAddressBase recv_buffer, - se::DeviceAddressBase send_buffer, size_t count, RankId peer, - const Executor& executor); - - static absl::StatusOr ToStream(const Executor& executor); - MoriCollectives* collectives_; // Parent MoriCollectives instance // This communicator's participant set (NOT the global MORI clique). `rank_` - // is this rank within the collective, `num_ranks_` the participant count, and - // `rank_to_pe_dev_` a device array mapping collective rank -> global MORI PE. + // is this rank within the collective, `num_ranks_` the participant count. int rank_ = 0; int num_ranks_ = 0; + // Owns this communicator's staging buffer + group counters (created in + // Create(), freed by the facade dtor before ShmemFinalize). Header-only + // facade. + std::unique_ptr<::mori::collective::CollectivesFacade> facade_; // Should all pending collectives cancel? std::shared_ptr cancel_; bool aborted_ = false; // Has Abort() been called? diff --git a/third_party/xla/xla/backends/gpu/collectives/mori_kernels.cu.cc b/third_party/xla/xla/backends/gpu/collectives/mori_kernels.cu.cc new file mode 100644 index 00000000000000..76b616a1942d22 --- /dev/null +++ b/third_party/xla/xla/backends/gpu/collectives/mori_kernels.cu.cc @@ -0,0 +1,19 @@ +/* Copyright 2026 The OpenXLA Authors. +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +==============================================================================*/ + +// This is the single device translation unit for the MORI XLA collectives. It +// is compiled as HIP and defines MORI_KERNELS_IMPL before including the facade, +// so the facade's device path (kernels + non-templated Run* definitions) is +// compiled here exactly once. The host mori_communicator.cc includes the same +// header without MORI_KERNELS_IMPL (decl-only) and links against these symbols. +#define MORI_KERNELS_IMPL +#include "xla/backends/gpu/collectives/mori_kernels.h" diff --git a/third_party/xla/xla/backends/gpu/collectives/mori_kernels.h b/third_party/xla/xla/backends/gpu/collectives/mori_kernels.h new file mode 100644 index 00000000000000..bff6bc1860d5e8 --- /dev/null +++ b/third_party/xla/xla/backends/gpu/collectives/mori_kernels.h @@ -0,0 +1,27 @@ +/* Copyright 2026 The OpenXLA Authors. +Licensed under the Apache License, Version 2.0 (the "License"); +you may not use this file except in compliance with the License. +You may obtain a copy of the License at + http://www.apache.org/licenses/LICENSE-2.0 +Unless required by applicable law or agreed to in writing, software +distributed under the License is distributed on an "AS IS" BASIS, +WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. +See the License for the specific language governing permissions and +limitations under the License. +==============================================================================*/ + +#ifndef XLA_BACKENDS_GPU_COLLECTIVES_MORI_KERNELS_H_ +#define XLA_BACKENDS_GPU_COLLECTIVES_MORI_KERNELS_H_ + +#include +#include + +// The CollectivesFacade owns the per-device staging + Run* entry points, which +// are non-templated and take mori::collective::DataType / ReduceOpKind enums. +// Host includers (mori_communicator.cc) see decl-only Run* methods; the device +// TU (mori_kernels.cu.cc, compiled as HIP with MORI_KERNELS_IMPL) pulls in the +// full device path and emits the definitions that resolve the host's +// references. +#include "xla/backends/gpu/collectives/mori_stub.h" + +#endif // XLA_BACKENDS_GPU_COLLECTIVES_MORI_KERNELS_H_ diff --git a/third_party/xla/xla/backends/gpu/collectives/mori_stub.h b/third_party/xla/xla/backends/gpu/collectives/mori_stub.h index 267cdf3a35fbe2..2a18155d71dc17 100644 --- a/third_party/xla/xla/backends/gpu/collectives/mori_stub.h +++ b/third_party/xla/xla/backends/gpu/collectives/mori_stub.h @@ -16,9 +16,14 @@ limitations under the License. #ifndef XLA_BACKENDS_GPU_COLLECTIVES_MORI_STUB_H_ #define XLA_BACKENDS_GPU_COLLECTIVES_MORI_STUB_H_ +#include + #include #include #include +#include +#include +#include // Inert stand-in for the subset of the MORI shmem host API used by the MORI // collectives/communicator backbone. These placeholders let the backbone @@ -72,4 +77,74 @@ inline void ShmemFree(void* /*ptr*/) {} } // namespace shmem } // namespace mori +namespace mori { +namespace collective { + +// Element type + reduction op enums mirror the real facade's non-templated API +// (mori/collective/collectives_facade.hpp), so the communicator's enum dispatch +// compiles against either the stub or the real facade. +enum class DataType { + F8E5M2, + F8E4M3FN, + F16, + BF16, + S8, + U8, + S32, + U32, + S64, + U64, + F32, + F64 +}; +enum class ReduceOpKind { SUM, PRODUCT, MIN, MAX }; + +// Inert stand-in for the real MORI CollectivesFacade. Header-only, all Run* are +// no-ops returning hipSuccess. Lets the collectives/communicator wiring compile +// and link without @roc_mori. +class CollectivesFacade { + CollectivesFacade() = default; + + public: + using AddressVector = std::vector>; + + CollectivesFacade(const CollectivesFacade&) = delete; + CollectivesFacade& operator=(const CollectivesFacade&) = delete; + + static std::unique_ptr Create(int /*myPe*/, int /*nPes*/, + size_t /*maxStagingBytes*/) { + return std::unique_ptr(new CollectivesFacade()); + } + ~CollectivesFacade() = default; + + hipError_t RunReduceScatter(const void*, void*, size_t, DataType, + ReduceOpKind, hipStream_t) { + return hipSuccess; + } + hipError_t RunAllReduce(const void*, void*, size_t, DataType, ReduceOpKind, + hipStream_t) { + return hipSuccess; + } + hipError_t RunAllGather(const void*, void*, size_t, hipStream_t) { + return hipSuccess; + } + hipError_t RunAllToAll(const AddressVector&, size_t, hipStream_t) { + return hipSuccess; + } + hipError_t RunBarrier(hipStream_t) { return hipSuccess; } + hipError_t RunSend(const void*, size_t, int, hipStream_t) { + return hipSuccess; + } + hipError_t RunRecv(void*, size_t, int, hipStream_t) { return hipSuccess; } + hipError_t RunCollectivePermute(const void*, void*, size_t, int, + const std::vector&, hipStream_t) { + return hipSuccess; + } + hipError_t RunQuiet(hipStream_t) { return hipSuccess; } + hipError_t RunFence() { return hipSuccess; } +}; + +} // namespace collective +} // namespace mori + #endif // XLA_BACKENDS_GPU_COLLECTIVES_MORI_STUB_H_ From 98d612580355a671ddeeef1d74c4351ddbef22ca Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Tue, 1 Sep 2026 12:10:47 -0700 Subject: [PATCH 29/43] Strip bazel-out prefix from generated #include directives. These generated includes cause build failures when path mapping (https://github.com/bazelbuild/bazel/discussions/22658); other strings in this file may also contain bazel-out and configuration mnemonics, but those should only lead to cache misses. PiperOrigin-RevId: 974634576 --- tensorflow/compiler/aot/tfcompile.bzl | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/tensorflow/compiler/aot/tfcompile.bzl b/tensorflow/compiler/aot/tfcompile.bzl index 911c917350d775..b6facb7eb12e4e 100644 --- a/tensorflow/compiler/aot/tfcompile.bzl +++ b/tensorflow/compiler/aot/tfcompile.bzl @@ -377,7 +377,8 @@ def _tf_library( sed_replace = ( "-e \"s|{{TFCOMPILE_HEADER}}|$(location " + header_file + ")|g\" " + "-e \"s|{{TFCOMPILE_CPP_CLASS}}|" + cpp_class + "|g\" " + - "-e \"s|{{TFCOMPILE_NAME}}|" + no_ns_name + "|g\" " + "-e \"s|{{TFCOMPILE_NAME}}|" + no_ns_name + "|g\" " + + "-e \"s!bazel-out/[^/]*/(bin|genfiles)/!!g\" " ) if gen_test: From e3ab2a2feb5799d1f2f8871fc0b83f1c285f29f8 Mon Sep 17 00:00:00 2001 From: "Patrick C. Toulme" Date: Tue, 1 Sep 2026 12:25:17 -0700 Subject: [PATCH 30/43] [XLA:MSA] Clamp a view use's prefetch deadline to the view's own time A view use extends the allocation end time through the view's transitive readers so the base buffer stays reserved while they read through the view. The prefetch deadline for such a use, however, must stay at the view instruction's own schedule position: the prefetched copy is materialized right before the view instruction itself. Using the extended reader time as the deadline lets the prefetch interval picker place the copy interval entirely after unrelated buffers have freed the heap, while the copy instructions actually run earlier, inside those buffers' live ranges. The result is two allocations booked on the same offsets and a verifier chunk overlap ("Value ... overlaps with another chunk"). Bound latest_prefetch_time by the use instruction's schedule time instead of the alias extended use_time. Adds a regression test in which the extended deadline places an illegal copy interval and the clamped deadline correctly keeps the base in default memory. PiperOrigin-RevId: 974642462 --- .../memory_space_assignment/algorithm.cc | 10 ++++- .../memory_space_assignment_test.cc | 45 +++++++++++++++++++ 2 files changed, 54 insertions(+), 1 deletion(-) diff --git a/third_party/xla/xla/service/memory_space_assignment/algorithm.cc b/third_party/xla/xla/service/memory_space_assignment/algorithm.cc index f01572da6645f8..08f5e7630a0b8d 100644 --- a/third_party/xla/xla/service/memory_space_assignment/algorithm.cc +++ b/third_party/xla/xla/service/memory_space_assignment/algorithm.cc @@ -6175,7 +6175,15 @@ AllocationRequest MsaAlgorithm::CreateAllocationRequest( // TODO(b/318886791): Rename boundary variables (here and other places) // like `latest_prefetch_time` and `earliest_prefetch_time` indicate // whether they are exclusive or inclusive boundaries. - int64_t latest_prefetch_time = use_time; + // A view use extends `use_time` through the view's transitive readers so + // the buffer stays reserved while they read through the view, but a + // prefetched copy is materialized right before the view instruction + // itself. Bound the prefetch deadline by the actual use time; otherwise + // the copy is reserved over (extended_start, extended_end) while it + // actually runs at the view's position, and the unreserved gap in between + // lets other buffers land on the same offsets (verifier chunk overlap). + int64_t latest_prefetch_time = + hlo_live_range_.instruction_schedule().at(hlo_use.instruction); // Control flow calls include kWhile, kCall, and kConditional opcodes. bool is_sequential_call = diff --git a/third_party/xla/xla/service/memory_space_assignment/memory_space_assignment_test.cc b/third_party/xla/xla/service/memory_space_assignment/memory_space_assignment_test.cc index cdd421527b137a..bfbc7d20331b0a 100644 --- a/third_party/xla/xla/service/memory_space_assignment/memory_space_assignment_test.cc +++ b/third_party/xla/xla/service/memory_space_assignment/memory_space_assignment_test.cc @@ -710,6 +710,51 @@ TEST_F(MemorySpaceAssignmentTest, ViewExtendedUseTimeWalksTransitiveReaders) { schedule.at(viewbc)); } +TEST_F(MemorySpaceAssignmentTest, ViewUsePrefetchDeadlineClampedToViewTime) { + // A view use extends the allocation end time through the view's transitive + // readers (`consumer` below), but a prefetched copy is materialized right + // before the view instruction itself. The prefetch deadline must therefore + // be the view's own time, not the extended reader time: with the extended + // deadline the picker may place the copy interval entirely after the + // c0/c1/c2 chain below has freed the heap, while the copy instructions + // actually run before `view`, inside the chain's live range, and the + // verifier fails with a chunk overlap. + absl::string_view hlo_string = R"hlo( + HloModule module, is_scheduled=true + + ENTRY entry { + p0 = f32[8]{0} parameter(0) + p1 = f32[32]{0} parameter(1) + c0 = f32[32]{0} negate(p1) + view = f32[8]{0:S(5)} custom-call(p0), custom_call_target="tpu_get_view" + viewbc = f32[8]{0:S(5)} bitcast(view) + c1 = f32[32]{0} negate(c0) + c2 = f32[32]{0} negate(c1) + cs = f32[8]{0} slice(c2), slice={[0:8]} + d0 = f32[8]{0} negate(cs) + d1 = f32[8]{0} negate(d0) + d2 = f32[8]{0} negate(d1) + consumer = f32[8]{0} add(viewbc, d2) + ROOT t = (f32[8]{0}, f32[8]{0}) tuple(consumer, d2) + } + )hlo"; + ASSERT_OK_AND_ASSIGN(std::unique_ptr module, + ParseAndReturnVerifiedModule(hlo_string)); + Options options = DefaultMemorySpaceOptions(); + options.dus_view_color = 5; + InstructionCountPrefetchIntervalPicker prefetch_interval_picker( + /*min_overlap_count=*/2, /*max_overlap_count=*/10); + ASSERT_OK(AssignMemorySpaceAndReturnStatus(module.get(), std::move(options), + /*buffer_interval_compare=*/{}, + &prefetch_interval_picker) + .status()); + // The c0/c1/c2 chain keeps the whole heap busy across `view`, so no legal + // prefetch window exists: the base must stay in default memory. + HloInstruction* view = FindInstruction(module.get(), "view"); + ASSERT_NE(view, nullptr); + EXPECT_THAT(view->operand(0), op::Parameter(0)); +} + TEST_F(MemorySpaceAssignmentTest, SyncDynamicSliceReplacementWithLateIndexOperand) { absl::string_view hlo_string = R"hlo( From 2b01994c7feaafc28a38c8f073ce2ee6e65f6083 Mon Sep 17 00:00:00 2001 From: Junwhan Ahn Date: Tue, 1 Sep 2026 12:26:31 -0700 Subject: [PATCH 31/43] [IFRT] Add RemapPlan constructor with input_devices_for_output_map and no mappings Adds a new constructor to RemapPlan that accepts input_specs, output_specs, and input_devices_for_output_map without mappings. Updates Validate() to support RemapPlan instances that have only input_devices_for_output_map and no mappings. Splits the old constructor into one with just mappings and one with both mappings and input_devices_for_output_map, deprecating the latter in favor of the new constructor. PiperOrigin-RevId: 974643048 --- third_party/xla/xla/python/ifrt/remap_plan.cc | 466 +++++++++++------- third_party/xla/xla/python/ifrt/remap_plan.h | 36 +- .../xla/xla/python/ifrt/remap_plan_test.cc | 199 +++++++- 3 files changed, 511 insertions(+), 190 deletions(-) diff --git a/third_party/xla/xla/python/ifrt/remap_plan.cc b/third_party/xla/xla/python/ifrt/remap_plan.cc index 4ea7e34451f3ce..72bcf73aae949f 100644 --- a/third_party/xla/xla/python/ifrt/remap_plan.cc +++ b/third_party/xla/xla/python/ifrt/remap_plan.cc @@ -168,6 +168,22 @@ int64_t GetNumberOfSteps(const RemapPlan::Interval& interval) { bool CheckOneInputForOneOutput(const xla::ifrt::RemapPlan& plan) { const auto& mappings = plan.mappings(); + if (mappings.empty()) { + for (const auto& [out_array, inputs] : + // NOLINTNEXTLINE(*-custom-deterministic-iteration-order) + plan.input_devices_for_output_map()) { + int first_in_array = -1; + for (const auto& input : inputs) { + if (first_in_array == -1) { + first_in_array = input.in_array; + } else if (first_in_array != input.in_array) { + return false; + } + } + } + return true; + } + absl::flat_hash_map output_to_input; for (const auto& mapping : mappings) { @@ -380,230 +396,314 @@ absl::Status RemapPlan::Validate() const { return InvalidArgument("Must have at least one input"); } - std::vector> in_used_buffers_list(num_inputs); - for (int i = 0; i < num_inputs; ++i) { - in_used_buffers_list[i].resize( - /*count=*/rep_->input_specs[i] - .sharding->devices() - ->AddressableDeviceList() - ->size(), - /*value=*/false); - } - const int num_outputs = rep_->output_specs.size(); - std::vector> out_assigned_devices_list( - num_outputs); - for (int i = 0; i < num_outputs; ++i) { - out_assigned_devices_list[i].resize( - /*n=*/rep_->output_specs[i] - .sharding->devices() - ->AddressableDeviceList() - ->size(), - /*v=*/nullptr); - } - if (rep_->mappings.empty()) { - return InvalidArgument("Must have at least one mapping"); + if (rep_->mappings.empty() && rep_->input_devices_for_output_map.empty()) { + return InvalidArgument( + "Must have at least one mapping or input_devices_for_output_map"); } + std::vector> in_used_buffers_list; + std::vector> out_assigned_devices_list; absl::flat_hash_map>> out_buffer_to_in_buffer_and_devices; - for (int64_t i = 0; i < rep_->mappings.size(); ++i) { - const RemapPlan::Mapping& mapping = rep_->mappings[i]; - absl::flat_hash_set* in_device_set = - rep_->input_devices_for_output_map.contains(mapping.out_array) - ? &out_buffer_to_in_buffer_and_devices[mapping.out_array] - [mapping.in_array] - : nullptr; - if (mapping.in_array < 0 || mapping.in_array >= num_inputs) { - return InvalidArgument( - "mappings[%d].in_array must be in [0, %d], but is %d", i, - num_inputs - 1, mapping.in_array); - } - if (mapping.out_array < 0 || mapping.out_array >= num_outputs) { - return InvalidArgument( - "mappings[%d].out_array must be in [0, %d], but is %d", i, - num_outputs - 1, mapping.out_array); - } - if (mapping.from.size() != mapping.to.size()) { - return InvalidArgument( - "mappings[%d].from and mappings[%d].to must have the same number of " - "intervals, but has %d and %d intervals", - i, i, mapping.from.size(), mapping.to.size()); - } - const ArraySpec& input_spec = rep_->input_specs[mapping.in_array]; - const ArraySpec& output_spec = rep_->output_specs[mapping.out_array]; - - if (input_spec.dtype != output_spec.dtype) { - return InvalidArgument( - "Input and output must have the same dtype: %v (input %d) vs. %v " - "(output %d)", - input_spec.dtype, mapping.in_array, output_spec.dtype, - mapping.out_array); + if (!rep_->mappings.empty()) { + in_used_buffers_list.resize(num_inputs); + for (int i = 0; i < num_inputs; ++i) { + in_used_buffers_list[i].resize( + /*count=*/rep_->input_specs[i] + .sharding->devices() + ->AddressableDeviceList() + ->size(), + /*value=*/false); } - const std::shared_ptr& in_layout = input_spec.layout; - const std::shared_ptr& out_layout = - output_spec.layout; - if (in_layout != out_layout && - (!in_layout || !out_layout || *in_layout != *out_layout)) { - return InvalidArgument( - "Input and output must have the same layout: %s (input %d) vs. %s " - "(output %d)", - in_layout != nullptr ? in_layout->ToString() : "", - mapping.in_array, - out_layout != nullptr ? out_layout->ToString() : "", - mapping.out_array); + out_assigned_devices_list.resize(num_outputs); + for (int i = 0; i < num_outputs; ++i) { + out_assigned_devices_list[i].resize( + /*n=*/rep_->output_specs[i] + .sharding->devices() + ->AddressableDeviceList() + ->size(), + /*v=*/nullptr); } + for (int64_t i = 0; i < rep_->mappings.size(); ++i) { + const RemapPlan::Mapping& mapping = rep_->mappings[i]; + absl::flat_hash_set* in_device_set = + rep_->input_devices_for_output_map.contains(mapping.out_array) + ? &out_buffer_to_in_buffer_and_devices[mapping.out_array] + [mapping.in_array] + : nullptr; + if (mapping.in_array < 0 || mapping.in_array >= num_inputs) { + return InvalidArgument( + "mappings[%d].in_array must be in [0, %d], but is %d", i, + num_inputs - 1, mapping.in_array); + } + if (mapping.out_array < 0 || mapping.out_array >= num_outputs) { + return InvalidArgument( + "mappings[%d].out_array must be in [0, %d], but is %d", i, + num_outputs - 1, mapping.out_array); + } + if (mapping.from.size() != mapping.to.size()) { + return InvalidArgument( + "mappings[%d].from and mappings[%d].to must have the same number " + "of intervals, but has %d and %d intervals", + i, i, mapping.from.size(), mapping.to.size()); + } - ABSL_ASSIGN_OR_RETURN(const auto input_shard_shapes, - ShardShapeVector::Create(input_spec)); - ABSL_ASSIGN_OR_RETURN(const auto output_shard_shapes, - ShardShapeVector::Create(output_spec)); - - std::vector& in_used_buffers = in_used_buffers_list[mapping.in_array]; - absl::Span in_devices = rep_->input_specs[mapping.in_array] - .sharding->devices() - ->AddressableDeviceList() - ->devices(); - absl::InlinedVector& out_assigned_devices = - out_assigned_devices_list[mapping.out_array]; - const int64_t in_shards_count = in_used_buffers.size(); - const int64_t out_shards_count = out_assigned_devices.size(); + const ArraySpec& input_spec = rep_->input_specs[mapping.in_array]; + const ArraySpec& output_spec = rep_->output_specs[mapping.out_array]; - for (int s = 0; s < mapping.from.size(); ++s) { - const RemapPlan::Interval& in_interval = mapping.from[s]; - const RemapPlan::Interval& out_interval = mapping.to[s]; + if (input_spec.dtype != output_spec.dtype) { + return InvalidArgument( + "Input and output must have the same dtype: %v (input %d) vs. %v " + "(output %d)", + input_spec.dtype, mapping.in_array, output_spec.dtype, + mapping.out_array); + } - ABSL_RETURN_IF_ERROR(CheckRange(in_shards_count, in_interval)); - ABSL_RETURN_IF_ERROR(CheckRange(out_shards_count, out_interval)); - if (GetNumberOfSteps(in_interval) != GetNumberOfSteps(out_interval)) { + const std::shared_ptr& in_layout = + input_spec.layout; + const std::shared_ptr& out_layout = + output_spec.layout; + if (in_layout != out_layout && + (!in_layout || !out_layout || *in_layout != *out_layout)) { return InvalidArgument( - "mappings[%d].from[%d] and mappings[%d].to[%d] must have the same " - "number of steps, but were %d and %d " - "(%s vs. %s)", - i, s, i, s, GetNumberOfSteps(in_interval), - GetNumberOfSteps(out_interval), in_interval.DebugString(), - out_interval.DebugString()); + "Input and output must have the same layout: %s (input %d) vs. %s " + "(output %d)", + in_layout != nullptr ? in_layout->ToString() : "", + mapping.in_array, + out_layout != nullptr ? out_layout->ToString() : "", + mapping.out_array); } - int64_t in_shard = in_interval.start; - int64_t out_shard = out_interval.start; - while (in_shard < in_interval.end) { - TF_RET_CHECK(in_shard >= 0 && in_shard < in_shards_count); - TF_RET_CHECK(out_shard >= 0 && out_shard < out_shards_count); - if (in_used_buffers[in_shard]) { + ABSL_ASSIGN_OR_RETURN(const auto input_shard_shapes, + ShardShapeVector::Create(input_spec)); + ABSL_ASSIGN_OR_RETURN(const auto output_shard_shapes, + ShardShapeVector::Create(output_spec)); + + std::vector& in_used_buffers = + in_used_buffers_list[mapping.in_array]; + absl::Span in_devices = rep_->input_specs[mapping.in_array] + .sharding->devices() + ->AddressableDeviceList() + ->devices(); + absl::InlinedVector& out_assigned_devices = + out_assigned_devices_list[mapping.out_array]; + const int64_t in_shards_count = in_used_buffers.size(); + const int64_t out_shards_count = out_assigned_devices.size(); + + for (int s = 0; s < mapping.from.size(); ++s) { + const RemapPlan::Interval& in_interval = mapping.from[s]; + const RemapPlan::Interval& out_interval = mapping.to[s]; + + ABSL_RETURN_IF_ERROR(CheckRange(in_shards_count, in_interval)); + ABSL_RETURN_IF_ERROR(CheckRange(out_shards_count, out_interval)); + if (GetNumberOfSteps(in_interval) != GetNumberOfSteps(out_interval)) { return InvalidArgument( - "Input array %d addressable shard %d is already used", - mapping.in_array, in_shard); + "mappings[%d].from[%d] and mappings[%d].to[%d] must have the " + "same number of steps, but were %d and %d (%s vs. %s)", + i, s, i, s, GetNumberOfSteps(in_interval), + GetNumberOfSteps(out_interval), in_interval.DebugString(), + out_interval.DebugString()); } - in_used_buffers[in_shard] = true; - if (in_device_set) { - if (!in_device_set->insert(in_devices[in_shard]).second) { + int64_t in_shard = in_interval.start; + int64_t out_shard = out_interval.start; + while (in_shard < in_interval.end) { + TF_RET_CHECK(in_shard >= 0 && in_shard < in_shards_count); + TF_RET_CHECK(out_shard >= 0 && out_shard < out_shards_count); + if (in_used_buffers[in_shard]) { return InvalidArgument( - "Input device %s used more than once in mappings from input " - "array %d to output array %d", - in_devices[in_shard]->DebugString(), mapping.in_array, - mapping.out_array); + "Input array %d addressable shard %d is already used", + mapping.in_array, in_shard); } + in_used_buffers[in_shard] = true; + + if (in_device_set) { + if (!in_device_set->insert(in_devices[in_shard]).second) { + return InvalidArgument( + "Input device %s used more than once in mappings from input " + "array %d to output array %d", + in_devices[in_shard]->DebugString(), mapping.in_array, + mapping.out_array); + } + } + if (out_assigned_devices[out_shard] != nullptr) { + return InvalidArgument( + "Output array %d addressable shard %d is already assigned", + mapping.out_array, out_shard); + } + out_assigned_devices[out_shard] = in_devices[in_shard]; + + if (input_shard_shapes.shard(in_shard) != + output_shard_shapes.shard(out_shard)) { + return InvalidArgument( + "Output array %d addressable shard %d has a different shard " + "shape from the corresponding input shard: %v -> %v", + mapping.out_array, out_shard, + input_shard_shapes.shard(in_shard), + output_shard_shapes.shard(out_shard)); + } + + in_shard += in_interval.step; + out_shard += out_interval.step; } - if (out_assigned_devices[out_shard] != nullptr) { - return InvalidArgument( - "Output array %d addressable shard %d is already assigned", - mapping.out_array, out_shard); - } - out_assigned_devices[out_shard] = in_devices[in_shard]; + } + } - if (input_shard_shapes.shard(in_shard) != - output_shard_shapes.shard(out_shard)) { + for (int i = 0; i < num_outputs; ++i) { + xla::ifrt::DeviceList* devices = + rep_->output_specs[i].sharding->devices()->AddressableDeviceList(); + for (int out_shard = 0; out_shard < devices->size(); ++out_shard) { + if (out_assigned_devices_list[i][out_shard] == nullptr) { return InvalidArgument( - "Output array %d addressable shard %d has a different shard " - "shape from the corresponding input shard: %v -> %v", - mapping.out_array, out_shard, input_shard_shapes.shard(in_shard), - output_shard_shapes.shard(out_shard)); + "Output array %d addressable shard %d is unassigned", i, + out_shard); } - - in_shard += in_interval.step; - out_shard += out_interval.step; + } + if (out_assigned_devices_list[i] != devices->devices()) { + return InvalidArgument( + "Output array %d addressable devices and sharding devices do not " + "match: Expected %v, but got [%s]", + i, *devices, + absl::StrJoin(out_assigned_devices_list[i], ", ", + [](std::string* s, Device* d) { + absl::StrAppend(s, d->ToString()); + })); } } } - for (const auto& [out_array, inputs] : rep_->input_devices_for_output_map) { - const auto out_it = out_buffer_to_in_buffer_and_devices.find(out_array); - if (out_it == out_buffer_to_in_buffer_and_devices.end()) { - return InvalidArgument( - "Output buffer index %d in `input_devices_for_output_map` but not in " - "`mappings`", - out_array); + if (!rep_->input_devices_for_output_map.empty()) { + if (num_outputs == 0) { + return InvalidArgument("Must have at least one output"); } - if (inputs.size() != out_it->second.size()) { + if (rep_->input_devices_for_output_map.size() != num_outputs) { return InvalidArgument( - "Output buffer index %d in `input_devices_for_output_map` has %d " - "inputs, but `mappings` reference %d inputs", - out_array, inputs.size(), out_it->second.size()); + "`input_devices_for_output_map` has %d outputs, but expected %d " + "outputs", + rep_->input_devices_for_output_map.size(), num_outputs); } - for (const InputDeviceRange& range : inputs) { - const auto in_it = out_it->second.find(range.in_array); - if (in_it == out_it->second.end()) { - return InvalidArgument( - "Output buffer index %d in `input_devices_for_output_map` " - "references input array %d that is not present in `mappings`", - out_array, range.in_array); - } - if (in_it->second.size() != - range.input_devices->AddressableDeviceList()->size()) { + std::vector> in_device_sets; + in_device_sets.reserve(num_inputs); + for (int i = 0; i < num_inputs; ++i) { + const xla::ifrt::DeviceList* in_devices = + rep_->input_specs[i].sharding->devices()->AddressableDeviceList(); + in_device_sets.push_back(absl::flat_hash_set( + in_devices->devices().begin(), in_devices->devices().end())); + } + // NOLINTNEXTLINE(*-custom-deterministic-iteration-order) + for (const auto& [out_array, inputs] : rep_->input_devices_for_output_map) { + if (out_array < 0 || out_array >= num_outputs) { return InvalidArgument( - "Output buffer index %d in `input_devices_for_output_map` " - "uses %d addressable devices from input array %d, but `mappings` " - "contains %d addressable devices", - out_array, range.input_devices->AddressableDeviceList()->size(), - range.in_array, in_it->second.size()); + "Output buffer index %d in `input_devices_for_output_map` is out " + "of range [0, %d]", + out_array, num_outputs - 1); } - for (const Device* const device : - range.input_devices->AddressableDeviceList()->devices()) { - if (!in_it->second.contains(device)) { + const ArraySpec& output_spec = rep_->output_specs[out_array]; + for (const InputDeviceRange& range : inputs) { + if (range.in_array < 0 || range.in_array >= num_inputs) { return InvalidArgument( - "Output buffer index %d in `input_devices_for_output_map` " - "references device %s from input array %d, but `mappings` does " - "not reference that device", - out_array, device->DebugString(), range.in_array); + "Input buffer index %d in `input_devices_for_output_map` is out " + "of range [0, %d]", + range.in_array, num_inputs - 1); + } + if (range.input_devices == nullptr) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` has " + "null input_devices for input array %d", + out_array, range.in_array); + } + const ArraySpec& input_spec = rep_->input_specs[range.in_array]; + + if (input_spec.dtype != output_spec.dtype) { + return InvalidArgument( + "Input and output must have the same dtype: %v (input %d) vs. %v " + "(output %d)", + input_spec.dtype, range.in_array, output_spec.dtype, out_array); + } + + const std::shared_ptr& in_layout = + input_spec.layout; + const std::shared_ptr& out_layout = + output_spec.layout; + if (in_layout != out_layout && + (!in_layout || !out_layout || *in_layout != *out_layout)) { + return InvalidArgument( + "Input and output must have the same layout: %s (input %d) vs. " + "%s (output %d)", + in_layout != nullptr ? in_layout->ToString() : "", + range.in_array, + out_layout != nullptr ? out_layout->ToString() : "", + out_array); + } + + const absl::flat_hash_set& in_device_set = + in_device_sets[range.in_array]; + for (Device* device : + range.input_devices->AddressableDeviceList()->devices()) { + if (!in_device_set.contains(device)) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` " + "references device %s from input array %d that is not in the " + "input array's addressable device list", + out_array, device->DebugString(), range.in_array); + } } } } } - if (!rep_->input_devices_for_output_map.empty() && - rep_->input_devices_for_output_map.size() != num_outputs) { - return InvalidArgument( - "`input_devices_for_output_map` has %d outputs, but expected %d " - "outputs", - rep_->input_devices_for_output_map.size(), num_outputs); - } - - for (int i = 0; i < num_outputs; ++i) { - xla::ifrt::DeviceList* devices = - rep_->output_specs[i].sharding->devices()->AddressableDeviceList(); - for (int out_shard = 0; out_shard < devices->size(); ++out_shard) { - if (out_assigned_devices_list[i][out_shard] == nullptr) { + if (!rep_->mappings.empty() && !rep_->input_devices_for_output_map.empty()) { + // NOLINTNEXTLINE(*-custom-deterministic-iteration-order) + for (const auto& [out_array, inputs] : rep_->input_devices_for_output_map) { + const auto out_it = out_buffer_to_in_buffer_and_devices.find(out_array); + if (out_it == out_buffer_to_in_buffer_and_devices.end()) { return InvalidArgument( - "Output array %d addressable shard %d is unassigned", i, out_shard); + "Output buffer index %d in `input_devices_for_output_map` but not " + "in `mappings`", + out_array); + } + if (inputs.size() != out_it->second.size()) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` has %d " + "inputs, but `mappings` reference %d inputs", + out_array, inputs.size(), out_it->second.size()); + } + for (const InputDeviceRange& range : inputs) { + const auto in_it = out_it->second.find(range.in_array); + if (in_it == out_it->second.end()) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` " + "references input array %d that is not present in `mappings`", + out_array, range.in_array); + } + if (in_it->second.size() != + range.input_devices->AddressableDeviceList()->size()) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` " + "uses %d addressable devices from input array %d, but `mappings` " + "contains %d addressable devices", + out_array, range.input_devices->AddressableDeviceList()->size(), + range.in_array, in_it->second.size()); + } + for (const Device* const device : + range.input_devices->AddressableDeviceList()->devices()) { + if (!in_it->second.contains(device)) { + return InvalidArgument( + "Output buffer index %d in `input_devices_for_output_map` " + "references device %s from input array %d, but `mappings` does " + "not reference that device", + out_array, device->DebugString(), range.in_array); + } + } } - } - if (out_assigned_devices_list[i] != devices->devices()) { - return InvalidArgument( - "Output array %d addressable devices and sharding devices do not " - "match: Expected %v, but got [%s]", - i, *devices, - absl::StrJoin(out_assigned_devices_list[i], ", ", - [](std::string* s, Device* d) { - absl::StrAppend(s, d->ToString()); - })); } } + return absl::OkStatus(); } @@ -652,6 +752,14 @@ absl::StatusOr RemapPlan::FromProto(Client* client, } } + if (mappings.empty()) { + return RemapPlan(std::move(input_specs), std::move(output_specs), + std::move(input_devices_for_output_map)); + } + if (input_devices_for_output_map.empty()) { + return RemapPlan(std::move(input_specs), std::move(output_specs), + std::move(mappings)); + } return RemapPlan(std::move(input_specs), std::move(output_specs), std::move(mappings), std::move(input_devices_for_output_map)); diff --git a/third_party/xla/xla/python/ifrt/remap_plan.h b/third_party/xla/xla/python/ifrt/remap_plan.h index d6c36040267bd0..475ea85bd47f3d 100644 --- a/third_party/xla/xla/python/ifrt/remap_plan.h +++ b/third_party/xla/xla/python/ifrt/remap_plan.h @@ -23,6 +23,7 @@ limitations under the License. #include #include +#include "absl/base/attributes.h" #include "absl/base/nullability.h" #include "absl/container/flat_hash_map.h" #include "absl/hash/hash.h" @@ -120,10 +121,27 @@ class RemapPlan { RemapPlan() : rep_(std::make_shared()) {} + RemapPlan(std::vector input_specs, + std::vector output_specs, + absl::flat_hash_map> + input_devices_for_output_map) + : rep_(std::make_shared(std::move(input_specs), + std::move(output_specs), + std::move(input_devices_for_output_map))) {} + + RemapPlan(std::vector input_specs, + std::vector output_specs, std::vector mappings) + : rep_(std::make_shared(std::move(input_specs), + std::move(output_specs), + std::move(mappings))) {} + + ABSL_DEPRECATED( + "Use the constructor that takes `input_devices_for_output_map` without " + "`mappings` instead.") RemapPlan(std::vector input_specs, std::vector output_specs, std::vector mappings, absl::flat_hash_map> - input_devices_for_output_map = {}) + input_devices_for_output_map) : rep_(std::make_shared(std::move(input_specs), std::move(output_specs), std::move(mappings), std::move(input_devices_for_output_map))) {} @@ -215,9 +233,9 @@ class RemapPlan { // and for each input array I a device list containing all of the devices // that hold shards coming from I. // - // Information must be consistent with the information in `mappings`, i.e., - // `input_devices_for_output_map` must duplicate, not replace, information - // in `mappings`. + // If `mappings` is not empty, information must be consistent with the + // information in `mappings`, i.e., `input_devices_for_output_map` must + // duplicate, not replace, information in `mappings`. // // Entries in `input_devices_for_output_map` are strictly optional, but // their presence may allow some implementations to be more efficient since @@ -235,13 +253,21 @@ class RemapPlan { Rep(std::vector input_specs, std::vector output_specs, std::vector mappings, absl::flat_hash_map> - input_devices_for_output_map) + input_devices_for_output_map = {}) : input_specs(std::move(input_specs)), output_specs(std::move(output_specs)), mappings(std::move(mappings)), input_devices_for_output_map( std::move(input_devices_for_output_map)) {} + Rep(std::vector input_specs, std::vector output_specs, + absl::flat_hash_map> + input_devices_for_output_map) + : input_specs(std::move(input_specs)), + output_specs(std::move(output_specs)), + input_devices_for_output_map( + std::move(input_devices_for_output_map)) {} + // `operator==` is more efficient with shallow copies. Rep(const Rep&) = delete; Rep& operator=(const Rep&) = delete; diff --git a/third_party/xla/xla/python/ifrt/remap_plan_test.cc b/third_party/xla/xla/python/ifrt/remap_plan_test.cc index a32fe7ffd867c0..6c572f80a5944c 100644 --- a/third_party/xla/xla/python/ifrt/remap_plan_test.cc +++ b/third_party/xla/xla/python/ifrt/remap_plan_test.cc @@ -97,7 +97,7 @@ TEST_P(RemapPlanTest, EmptyMappings) { /*shape=*/Shape({2, 3}), /*shard_shape=*/Shape({2, 3}))}); RemapPlan plan(std::move(input_specs), /*output_specs=*/{}, - /*mappings=*/{}); + /*mappings=*/std::vector{}); EXPECT_THAT( plan.Validate(), absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, @@ -693,7 +693,7 @@ TEST_P(RemapPlanTest, InvalidInputDevicesForOutputMap) { std::move(input_devices_for_output_map)); EXPECT_THAT(plan.Validate(), absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, - HasSubstr("not in `mappings`"))); + HasSubstr("Output buffer index 1"))); } { @@ -719,14 +719,28 @@ TEST_P(RemapPlanTest, InvalidInputDevicesForOutputMap) { std::move(input_devices_for_output_map)); EXPECT_THAT(plan.Validate(), absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, - HasSubstr("references input array 3"))); + HasSubstr("Input buffer index 3"))); } { + std::vector three_input_specs = {dummy_spec, dummy_spec, + dummy_spec}; absl::flat_hash_map> input_devices_for_output_map; input_devices_for_output_map.insert( - {0, {{0, GetDevices({1})}, {4, dummy_spec.sharding->devices()}}}); + {0, {{0, GetDevices({0})}, {2, dummy_spec.sharding->devices()}}}); + RemapPlan plan(three_input_specs, output_specs, mappings, + std::move(input_devices_for_output_map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, + HasSubstr("not present in `mappings`"))); + } + + { + absl::flat_hash_map> + input_devices_for_output_map; + input_devices_for_output_map.insert( + {0, {{0, GetDevices({1})}, {1, dummy_spec.sharding->devices()}}}); RemapPlan plan(input_specs, output_specs, mappings, std::move(input_devices_for_output_map)); EXPECT_THAT( @@ -764,6 +778,139 @@ TEST_P(RemapPlanTest, InvalidInputDevicesForOutputMap) { .status()); } +TEST_P(RemapPlanTest, InputDevicesForOutputMapWithoutMappings) { + ArraySpec dummy_spec = GetDummySpec(); + + std::vector input_specs = {dummy_spec, dummy_spec}; + std::vector output_specs = {dummy_spec}; + + absl::flat_hash_map> + input_devices_for_output_map; + input_devices_for_output_map.insert( + {0, + {{/*in_array=*/0, GetDevices({0})}, {/*in_array=*/1, GetDevices({1})}}}); + + RemapPlan plan(input_specs, output_specs, input_devices_for_output_map); + EXPECT_TRUE(plan.mappings().empty()); + EXPECT_EQ(plan.input_devices_for_output_map().size(), 1); + EXPECT_OK(plan.Validate()); +} + +TEST_P(RemapPlanTest, InvalidInputDevicesForOutputMapWithoutMappings) { + ArraySpec dummy_spec = GetDummySpec(); + + std::vector input_specs = {dummy_spec, dummy_spec}; + std::vector output_specs = {dummy_spec}; + + { + absl::flat_hash_map> map; + map.insert({1, {{0, GetDevices({0})}}}); + RemapPlan plan(input_specs, output_specs, std::move(map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, + HasSubstr("Output buffer index 1"))); + } + + { + absl::flat_hash_map> map; + map.insert({0, {{2, GetDevices({0})}}}); + RemapPlan plan(input_specs, output_specs, std::move(map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, + HasSubstr("Input buffer index 2"))); + } + + { + absl::flat_hash_map> map; + map.insert( + {0, + {RemapPlan::InputDeviceRange{/*in_array=*/0, + /*input_devices=*/DeviceListRef()}}}); + RemapPlan plan(input_specs, output_specs, std::move(map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs(absl::StatusCode::kInvalidArgument, + HasSubstr("null input_devices"))); + } + + { + ArraySpec f32_spec{ + /*dtype=*/DType(DType::kF32), + /*shape=*/Shape({4, 3}), + /*sharding=*/ + ConcreteEvenSharding::Create(GetDevices({0, 1}), MemoryKind(), + /*shape=*/Shape({4, 3}), + /*shard_shape=*/Shape({2, 3}))}; + std::vector in_specs = {f32_spec}; + absl::flat_hash_map> map; + map.insert({0, {{0, GetDevices({0})}}}); + RemapPlan plan(in_specs, output_specs, std::move(map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs( + absl::StatusCode::kInvalidArgument, + HasSubstr("Input and output must have the same dtype"))); + } + + { + ArraySpec layout_spec{ + /*dtype=*/DType(DType::kS32), + /*shape=*/Shape({4, 3}), + /*sharding=*/ + ConcreteEvenSharding::Create(GetDevices({0, 1}), MemoryKind(), + /*shape=*/Shape({4, 3}), + /*shard_shape=*/Shape({2, 3})), + /*layout=*/ + std::make_shared( + xla::LayoutUtil::MakeAscendingLayout(2))}; + std::vector in_specs = {layout_spec}; + absl::flat_hash_map> map; + map.insert({0, {{0, GetDevices({0})}}}); + RemapPlan plan(in_specs, output_specs, std::move(map)); + EXPECT_THAT(plan.Validate(), + absl_testing::StatusIs( + absl::StatusCode::kInvalidArgument, + HasSubstr("Input and output must have the same layout"))); + } + + { + absl::flat_hash_map> map; + map.insert({0, {{0, GetDevices({2})}}}); + RemapPlan plan(input_specs, output_specs, std::move(map)); + EXPECT_THAT( + plan.Validate(), + absl_testing::StatusIs( + absl::StatusCode::kInvalidArgument, + HasSubstr("not in the input array's addressable device list"))); + } +} + +TEST_P(RemapPlanTest, CheckArrayCopySemanticsWithoutMappings) { + ArraySpec dummy_spec = GetDummySpec(); + + { + absl::flat_hash_map> map; + map.insert({0, {{0, dummy_spec.sharding->devices()}}}); + RemapPlan plan({dummy_spec}, {dummy_spec}, std::move(map)); + TF_EXPECT_OK(plan.CheckArrayCopySemantics( + xla::ifrt::ArrayCopySemantics::kReuseInput)); + TF_EXPECT_OK(plan.CheckArrayCopySemantics( + xla::ifrt::ArrayCopySemantics::kDonateInput)); + } + + { + absl::flat_hash_map> map; + map.insert({0, {{0, GetDevices({0})}, {1, GetDevices({1})}}}); + RemapPlan plan({dummy_spec, dummy_spec}, {dummy_spec}, std::move(map)); + EXPECT_THAT(plan.CheckArrayCopySemantics( + xla::ifrt::ArrayCopySemantics::kReuseInput), + absl_testing::StatusIs( + absl::StatusCode::kInvalidArgument, + HasSubstr("kDonateInput is required if multiple inputs are " + "mapped to one output"))); + TF_EXPECT_OK(plan.CheckArrayCopySemantics( + xla::ifrt::ArrayCopySemantics::kDonateInput)); + } +} + TEST_P(RemapPlanTest, Hash) { std::vector plans; plans.push_back(RemapPlan()); @@ -784,8 +931,8 @@ TEST_P(RemapPlanTest, Hash) { /*shape=*/Shape({2, 3}), /*shard_shape=*/Shape({2, 3}))}); - plans.push_back( - RemapPlan(input_specs, /*output_specs=*/{}, /*mappings=*/{})); + plans.push_back(RemapPlan(input_specs, /*output_specs=*/{}, + /*mappings=*/std::vector{})); } { ArraySpec array_spec_s32{ @@ -999,6 +1146,46 @@ TEST_P(RemapPlanSerDesTest, ToFromProto) { } } +TEST_P(RemapPlanSerDesTest, RoundTripWithoutMappings) { + const Shape shape({4, 4}); + const Shape shard_shape({2, 2}); + const DeviceListRef devices = GetDevices({0, 1, 2, 3}); + + std::vector input_specs; + input_specs.push_back( + ArraySpec{/*dtype=*/DType(DType::kF32), + /*shape=*/shape, + /*sharding=*/ + ConcreteEvenSharding::Create(devices, MemoryKind(), + /*shape=*/shape, + /*shard_shape=*/shard_shape)}); + + std::vector output_specs; + output_specs.push_back( + ArraySpec{/*dtype=*/DType(DType::kF32), + /*shape=*/shape, + /*sharding=*/ + ConcreteEvenSharding::Create(devices, MemoryKind(), + /*shape=*/shape, + /*shard_shape=*/shard_shape)}); + + absl::flat_hash_map> + input_devices_for_output_map; + input_devices_for_output_map.insert({0, {{0, devices}}}); + + RemapPlan plan(std::move(input_specs), std::move(output_specs), + std::move(input_devices_for_output_map)); + + ASSERT_OK_AND_ASSIGN(RemapPlanProto plan_proto, plan.ToProto(version())); + ASSERT_OK_AND_ASSIGN(RemapPlan plan_copy, + RemapPlan::FromProto(client(), plan_proto)); + + EXPECT_TRUE(plan_copy.mappings().empty()); + ASSERT_EQ(plan.input_devices_for_output_map().size(), + plan_copy.input_devices_for_output_map().size()); + EXPECT_EQ(plan, plan_copy); +} + INSTANTIATE_TEST_SUITE_P( SerDesVersion_NumDevices, RemapPlanSerDesTest, testing::Combine(testing::ValuesIn(test_util::AllSupportedSerDesVersions()), From 822a287a673dc54091b6d32d72254661e988a962 Mon Sep 17 00:00:00 2001 From: Alexandru Paiu Date: Tue, 1 Sep 2026 12:31:35 -0700 Subject: [PATCH 32/43] This is an automatic update to a device compatibility allowlist. PiperOrigin-RevId: 974645403 --- .../compatibility/gpu_compatibility.bin | Bin 120792 -> 124380 bytes 1 file changed, 0 insertions(+), 0 deletions(-) diff --git a/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin b/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin index 98c198ee6baf629972c8cdc7a07069448df12fd6..031f8772277612e2f2858aa9ae7ac03443099374 100644 GIT binary patch literal 124380 zcmaI94_sVVU*~^8%)?XJh!HDRtXO#>PnXRO!vI6XE@rV}#EKEK7%{`2fsk|{^T0IE zw6a;OY!@S2F=E7u6(d&6vJopr%wiWKR;(DYV#SCRD|Y#H*(}>-E1PBO@BN*7Cb{R! zUHf`{JD2zU+|N1pd(XY++%tD>avb^Zkdy6X@qdT^JI&WU9d$Qoue_%1A zWz*(I;KjFXmJQ1n9g)ZKp7Oe2yz-RSS@`;I*rLao!ZzOhz78K#{$sEqbO~aT z$<;Y{5RRg>;Eb|_L|KGVdLnA2rd<- z?iFG2o?u;Vy?Z{V+vz99Tl8M{-2<& z{SDiH8x}%WNIc4Od8K$~#01OR%Ig#3HKe>gf;XsoZNO^Kl@Qx3?<=npPnqMdyei;X%Buvn z%e8!i2(zq@^{d8fM_unInw|TP2-NzQ?bp&3lys}~W zq9gKH9#mczj8~rWItyR_HM`tK*z#Y}{}U}Nk14NP#%qc7-hI3QA5mTtumN;2Vw&Y? z<(2t4cOUze*C%LemDe^bgsza-V|hV&J!iZE%IgGtZ^o``7`Anf{-3C2xlr|kuNkjR z)_eEyRrs{>8ib9b8zM$nuKzQ8n-7iGnDV*@Z&6+=um*J1#5&6%Rj=StH(&M2>nJ== zc@@C2(Cz#MV+WS&l-FhBXJy<@y)S?|r)TkwA6H4BTO zYa$j{URPd^f0Mh9UCQe*ctCkMutIb>5-caxHf_eMK=IS?&6{>z>tU;F)QPyn@)n=6 zFLGZ|eVFn3KVckCw9t93aPJ`N2GH&PvF-CFY{9gFZ$^_uTSw%Z`;9s7HHL^LqS)Mb zEHmgCBcen=20+9Lo!l4^B?81g_a4i{I1wYNiELt(3QiJnqMFDd*0>)TBU*?8Vw1`) z5F>=#qZAXn+@CBE!$c!dNbGQ*vTW{IqI}j6IRvS2CW!%}fhZ()sOGHfA5l$U>D(;- z^dAb$GLjUVg!TVQ$E|(!PeMjk@!koUNBHvMMbapY&QU#V4LQ<84icS zx%B*tkyxP0x#gpsh9%Gj(%N=15DYpW_-NN)b!fAUwyy_mGUyzBnz6;S-QKcb1!%L0 zW3o&=LGCu(txwhGeINZQ ztR8(fddYvho%#f;oX7qdV`0^YIj};sImFW}>$YDAb67Y3j(hFtzL>ZeIagfy`A9t0 z`XcS$XPZXYGGqE7wqY8{d9=;Nc;b9FqugYow5mK9jeXvE-Dmqz*f87I*=1Iw%Y4^o zn3MQk3Iz3{uA0iddY8U z|2it13qINw*aF)Aw0CPSZKJHUo9$ln(T~7J(AT4<%i>;G^ln>v$49>aYe8R-u8XU8 zKJn3R!$N3xmb~7jT@>RS{pSbsQUJ?BzhLXzM1R^xUk_VdVGM^}@)CQoHP!B3_R%I` zQ)r`JE!X{Sq95?+D^mWBbInIT3>!t?fL`Vo zq(112sZXlyzT>ma0<48?3fM++exXltQ|5f)quqvu(C)mQ;V`6eE?j&uQp)F3r3Z79 z56ecoh*sL+g=kMO)iJYt%12)ZTcfXCBqTSCFk+Edd3m7RabEV(Ct*`)M+wQj(FX#- zxbvouehQXE8zmTKIp>)ky62-?fi)nQHm&V;+xA%>Z6j_zUK--+w#<*4oLe4E8 z?KCWbHej@E&y!-#nGbxlYp^=BSw`F1>-GVM%MRutOBmXVG~Qjm<6g_kjs=`Ie01ZmL3A}n z7esg0N4E%TMwf4NrN>I0HCVoCACFbK!@R^I(`@YTG0u8R~CpngEH_Sdwek1jv z;~5`eABeRdXMUP+@XONO&#O!x6Z@>w$^f}7AL5}h$ zrxMLBtdshE#QemXRbAA?!Fg#&P9^#vGei1jtGOK&qo;2UV0(+d}|N7 zQG5pZ^w#%n*z_H{ym^?e?~!kI>l>pE-t7;;mubIVw$JuO=yD0U?#9n28MKv_N{=2% z#vHnXbe-cDzu3`xsW;Ld>Fek{-_aLoPbT7#P&iUrUREJ_7HvIcZ1CwV<1%cQ>yj|J zZc1KSh$LLs{Vt4E88fQw+=17A!QP)aSQk1OFQnQ-d1+~9yuB<$dsy}He!ANA7QMHA zPr|q2Hp_$MpxdK;q}y6Fo+vNFTlzHfU3W0Rj)(cX7X1V#nb-FcY=6P^6&@pUhy@}a zZIbl|_8-RT%o)`#@4#zS9p_*&ekl4+58GN_RVXOU$6599e)?Nn??&bQBz)_l+J9IM zx~1Eecd8u+Sf}g#it*l4?eiu0u=0+>X3)hf@5+?-u#fi##(P|OzXuN}?`2pcx?I~k zm6Lzu{9`K5KL7Hb;n;D|yvOrdSjm5O^CWFxm2*kvfcuGJ{@x@SXy9Rvt@QiDd~OmA z=wwcKgoxqMz&fw*9T+HM-kc+Bl3ITYuE$o>umrjm0z0RSzA~Fnt}~rg<~p;2MjxM} z|H3V!NtN*ge7oB|_PMZJbkcT>H|?sDS!Ty6VV(AV*?6bhdb<9X;3LXA2Ajoq!txHL zyoY?e-#6Z(_vY+fcu09K!J5$3+TJfoy*GWlkJq?45xv*@7Dz1RCgc%kxMht;ABSl+=eO1<}d zyr274Hz%U^dY^zVK5w%yZ0jw?e1s<_Dep$s>73j!-lF$7& z==urCVQOD1eCou#^*q`gE;!!u@52Uv%k`9dGKSes0*&;!^DMUzyJ&PBPu9A145>Ok z17GU0>llJ<)9G45FJkc)S@HDVad%t14MenWaYw(!z9)XRZ>nEiCmxAS`<$O;0 z_>dd*LMe| zkI%EuruH{8tm<_fzWQtS_5s)q?Q7cbSVeiOq*3lc(u3r-^dPw{*%1jw%BY^x#dh1O zt-c0ZGxuZZw#qQ?VLQ!Xc>;~D?Eja)mOFbs zzR!KT>l;wMC*YeO+UHpqwl&A~pGdbSZiSr&xUTOT#&->^)&Akb%69}dg07#C`u27w z8QE7jQ$D_*8s7os`!PJEe79g>bVWqEzS5sKbqDKv8pfK8V&!`hzW#T19%^74w6SI4 zVSSnFmHR30GhLNT7{$;|nl>Qgw=P2RaUH&;+Ub~Rf5Xn#eHikL5vBbAUKF*pYp^mgIWg?r6WGq@)QIYCT zn|-|BG~S|@^3!ei4S0+49)}I0i`(9*vC*85_u-Rny+!Z!&V*+v?>)9JMwe%Mr^a-j zqh9}yd-W=PuU><99KE-^Ua7~g+jZ!IZM|l*w?xPN-hw*LZ^9efZ0!^*i7xaXGn^M6 z)$e={-?NmMse?%Av9e=jrQt|%3K#c&K61YIe4l)BAv)LF8t*#a-WMrujm9G95<$Ki z4I~5n7piCtfBy3&%=yTdHW%X^J(0ePFZA+sMR!LZ->mnZztqti36>vY0#fetL*L0< zy=lkk)dO+MejZ4o$!UfvFqFn zTcjP#5SI5j?*Dzf?;3B>d%f?#qsn^@)`hN-kn7p`U`W1dJ8<2RYyBqL8ItSueEhqZ z{{jk=2(f;JPdQ#)L_X{FzW2abm65CJ*9M=A*mZ7#EzurkiLYQIGq)5-%`J6aj74Gn zK4rXX%9u87iE5?l_cpvnl`#*CqH7>L^H**e%wMhh_y?bJ_n}_-ABAt6wfCU_mW6JI zzFy{Ty!*hPPBrUwj;U(WA%P+ez(>2`1zKJ@c;y%u58 zmh-9Ckp}np#?*F);DxGQ`)pT)E|+*%uil$Ow!Ax?Ac$2(z4o6My+u*remxW6# z5nY7jJ|9h#^>r-Ey7YeB7BBu>{8f8jZo-?IZ8iy$XF$6M@l4GtS5$Pc=q&hnKmI+g zx9GiPJ_g_a0h>9nLUcK{cbj+@uuj*z*LaKG>)i=oShx4F8Ma7UoFP2*mSLik^zpuH zyhZQzz5|ac?>Sf(x<*3AUF0NL&T}63;p1xiO<2>k?Ror^o1qswJEN-&M%-D$}A5^|$ zupx8_LUJtgAQk1M&Ww-mq3?6+n^e9J(3U9Q9asQef#qBEkA84W=o z3C5-w+2%Y>pMw|R*v=dLwuX>Ar0UCcsSd95@VfEcKzA1SB4pph zmk*}_uJiDQ@m(|8bbELWKCFC4U?b@I38}B#*H!X4<>ULQ@f}dUAHze+cMBFqS7dyd zR}RQ^-T8ro^*s$^O-8ZuJqgcY(Ca;xHLwjPwU&wW7`mM6TL;&PW*_f2jkoB%K5xKh ze$r;+ut9Wj?ah6I9-s633TGL8BiC@R@29ZDFSx#Ptyv@aNurBo8JlPGIY(qSAI!s9 z7^^Z0Rb5ZPSO3H=rw+FEi#!uW$a&#@Z&B*B`1sy3zDsDmb-e-aSH2Ul0dz4UJtnG( zCpx76nD_BM@`G-jMep@K1TR$H`)pr?E|++CE{DsVJhZwEUogIT%J(dMX~oV(BW(FP z&y*1lb5U01#C&}37~grc-um8x$Cd9iEP<|tFl{XujkgC{X%9<2-cO!&b0K=K_YruW z^3I0ki%xr!XZ_5FoR4Aj8+>|wFTk?Ry(!J`;T&reuJd^th9q;*w9WK2;U+x!$F}bz ztRG#NvC?PRC{e*KM|v-Qk>v#+@5dvqcbD>h3@*>SdF$-J3el~Px!&>eN8Vq9e|iqW z^PNU_B*agYSy#=r+nhk+C;Pj>!ItG-DgSjh=U>I&V~nYCSQjN4I0?tOux%l%=5vj8 zgDfv%&_sku#2ztD44_>?Bku#r#ixtdMKjOm9A)U79&T}ucYMM=-kI>M@3Pq*+ZUtD zBiIEQEA&;BmA0|ynwb#-e!%kW!k|U@yOeKULfw{{kHKIz1RCT zJoZQS@tT3f(KQp&KlEKZm%5Kx^6`DL)y+qX@;w62Q@+`-e9;m4_{uZEy%%Gxk>1vq zdd)X&{QR$+RZa-4&d(K7)~-43>HJ)VPjQp$-LC{}8r>Kn*W{jLWm(w;#(zG(9~j?p z<@+AILHRDj8qw7dlAE4*pd4Fg)5rICo12?j<$Dai_lI_Ui(vad$NfL?IDU3+DgthL zc>WEY&dYUE&N|zAb8rnlqRJVDjiMVMI8~hPwyxCk1Cu_!4~*}i^1TlaE8k66HM(L# z^4)haQeN80^QJz&r`p}Tlqlb4;T!MRd8vkNa-yyf>FaT@_xbk9bDYbp)48~5yhZQL z#dY|g@*aZ?p-T{wi`4yZCHKERzK5c2os-J<0ooGfy8{cLD=@yP@hfwC*?8#so;ALO z%J&p}^$+Y^)WOz%mj0iRwkPAZQa)RJd~X@wCA8i=+<^Bh-wD_Nx)|{=59K`n@8f&8 z!_7ln`DVh4l3|^po9atf{9767)Tl-$9+Ki z=;ZrfyuJB(Q?x&3vnf~-UFgRjomYWSIQ6}h=iOT}G<8;K-~w_?~H$8?jB?*jW??CIb)YF&{;S!zsF$9kQ+*Iob24b>K|!uyr~ z5Nr%xBV4XST`#uvo{vVt6@h>=>Er#tcn>P?`|v{Ly$P#Ew>Rwm-slVbKJ>Gc;kFxo zVajs$^XxO9-s5>37QX7{f##GN%eL?tg6n;~0%KL?u9_dX44=PX=Q07CM%P7sQrFyG zez6fL4MccF!HAFdhsJwMdEbM#DDM?m1G*f%WnaD+Pwk%PJW0K;OKxABeYCk;$V4am z*G0&3Fvj;qs;pwzWR0Dl%P{0zuSxWvo6>i^BeM;sH-e*jJWGRgMplXGUo2* zd|V+GS(fuRkL3xpI@h1TkYo(0`I3*|)xTrcaRXL^u7ogsMWDSm9%?_&{`+{J>UORbZbrTuIWUyYwV4-k*>19YZHLQW$( zotqn`oK@8~zYd?{!tbsBC~O#AlCb6*I-=+Jx!}Y3hIyZIj`X?v6I10Jf;Xsg_SwD& zU9K)CHRtoFoFcZ>`}2}1Crg#n1K;`~dw-&^dHUli!n7lvv3oI4MLQbs@qW*Ei{6`; zci_XydjZyhuKtm?+eYbw)_uH#3Df?S_fdF2c^AO4(Csi5mus(kjVyISaGj4U#&=hZ z+b_e{Ub6Qi0h>lQW_&w%-$GBa%IDf!?o6OvWsoF!kZWxE`;GVDv!*PmTSC^!wTjP8 zKIQtxB`@`!>u>C8{e-Uy%rQXfJXEl7a+tyw+zT0TM z92|O=0;fYtY=x~^=FAqw#lVZ@=ospyqKa*{r0s5wt`NcbE;$e zBxUM0dDql`M&3j_c@nzN42ie0+~4-Mlm@-zVVPAv?#puv~Qe z|L)=3eKJ_){)W4Rb-KkX*!){AMD|_0PvQ_@bMOdgXgm_;1)bE`VjB+o5g9cVN5|;C#>dWF%P; zPQ5g^n)Ny#SB?L+X-^#TR9;?zPt4l>{jf=NBZSy##{8>Tz>$s zRo-i|J-QGfeNNgt{5M^+A3wPN;Um6C8Z+=i`6d_%E`JH!nBgQ++m@g!QBABEAAY zd)!qX=z}f!l<{Q0yALg@j3e+owGY{_e9;k)+G2aRw8bLU>Ers6@fN+e{yp%mUYkW> z^FPA)hcM%xs`7YyYxJDs4ET7zXS_x4^?nCFth^Us@{W;uk9VNAHCToBx{r78U%NRG zz1RCFJfOS_U|HyPn9GoMp8BqY`KS6Md)$~x@&nRHc-ZA9$I~TFBsoE<$D&schUB3ge`LuJx7SId>1CZF(2PM#&;gA zw>{p1PblAMSOQ%Okv@-7ZJqmwB_Ho62h91eypOseErjH04zBBU_=i*b zn;BR0DVgw`FWU9mWBX!sn~Wu-UWwj|=b{&TX$uE^JB@F?@;w8eRb!C`*b4Iq!*IEd z_F>#}uB|o7vb2THXoU<-oFuwE1_$0_egn2-?!iPW=WQ2}kB%XF>bdwCG^BmhRxMQVwYcx_pCzXwd_C0HS3zhFFc#iU|gRL=NFp96#J<-}0 zkuL^Xe0*;i-z7D_a05Q;*!7)&4WO&WH#LsIw=(j@^PRN$c^~g1KjP*hro0cqo0RuH z+ZUnRNIdc`k34@qN{5<{Ugx97c;_nbHuyeQOz*L5f-TW*2l19Ue#u8`rE8we-|lx0 zWvTJaAX=G!^KO3&Hm~mSreV5HnS<{B_N%##Ptew$vU9l&3!&RFNJB^-a8|?lIuYiAOL#qYmXe0+}%xw&XkzE8lbm2WOA7u`Pnr5R&X zcE)>B&oh^>PUqrf<1Kn`E-t}0I&Bt%&7zwi+_qd6kiKxx$M>G`om9Se;X}%I5!Q^Z z&f_Z+=`xq|=v`&@d+}tvvCno)Z(8*}_M==kxn|40NaV3z<`QR#CL+Lkor70ktjgF^ zV}Y08qiQ@9hs~fHC%8;F3Em+%=;L|Mcup$MyYL$2xd>}US4Wt<%Ii6ubsyj06}PVS z%J(RItJ^+a1+Xl1JEl*+6by#selP^rIk;kcchP#=!DYC94^{#;jc$yPbD}j7=`Lj+ zNb>CMhxpT(Ks$?}czUn*?+I6BEyLpK-ubB?bN9DdX&;BzMD4of!1B;JLIHH9lodFNuW#9TY~xhgOBf<#&=mAr#Ij&%6A+#h_2>QojcNXo@1R1QoL>L@V|HK zoNfcsUW$$8kgQYw`*;-Lzi}bMIV$Ua!S*YFWwCyN_1XAv?QShEuV8e@_kXM#G~QB9 zjaj}bKICm%%_Y17ANYDZZ~d@Ibn*Oo z;L%r8b;=OEH(yV{tCe>yEEnBA_xY)L<@QKe(%hd0r2l2@w*^)nLOLDI#%eUd;o#s+T z12GF<;+?(1)7O*7M%-hZP;(rgqAh;i*6zYe(B%^j+f6-f=i{{TDpp=6;p;zVd)2@; zm`hkDF0hXtgB{O#?V~)oKbvq8~zB! zrWvDZ-r{}HsyeO0>d}P>>0eUw)6Cd!`}m&t3D>t;`BuPp#_fGAf$h@wZV+-Tx;kDg z<2(D*c~{3eUFX-0x9Hi$hsWwFeDMcuHUt|(H%Q3+VppO)@adlOcTu9%Q~M0;LLec!t8<9q6=n~xIZ`z(CpXYG7c!!|kID}+1l2&Jz3jc}cd z8^(84jWJ$_&o$foFbW$+mn6Q5`aYafG;^1I?ko45d7pBQ{A8;BnK4z)A$Xyx|32Fn zq1#{_VCtW`x5`7S>wm%Z&B#-}XW@H4VAsD9woF5wBgFSSUvkM@NX*Cgj`5vG>phOQ z;1kMs8kRuU^3XS!noC;r@qJ>{J&sZ3dl+7)e6wH$=(3IPrC=FP)^g5@mQQcJdW~<6 z^6iB0oU!+z87A+|oFVe?(Z|soMrY8coO`C6Nw)Qte-}Qa%2|XpqpKsNeRuFr@yhcA z>ps50pK|k3uY8Zf1Io7mmW6JI{@HCGyj%+Z5M1Zwit*h=>#gr)_}cf|ECHKFH%6qd zPZgy;&z^BVH-UDOYomBd-yr|ank>HuA2;`I+5A1p$|RpPd`hm@xsDbR3oJ8)bGFdy zIzMyGJy&W}orCZqRp(;Z9{u(#;{_>?u~1)Su=I;8I<>6R-me+&O}6po<5l=llg$QU zzFCVadZ_#_bufu2mmCeRrL+BEgcSXv3*2nv?F}L2L_j-Sdwn=&K z!b;HPTi)((|8nus_3m`NGez(9J_FxBZL?lNY3@}Tr+7$Pt})AR zyZ#w-YA)h7ygtwNpMk~EH4~q@+hq=@?LLM_RlT-gVRXA^t$MXl1?Q=sVJ<~^kpS`{xHp#jX;&XWXzuvj0->Wmr zcJgmFrt{ck>hKYK0x#j5@~O`aSnL~YR`HsW{baTS(f*Vl&~zItghbhdlS6~h3stM^^ z_zi1nh`QzD`^>nT4?PbOgzx@@UFTxh9>;!@NcUNlrBzN1T<7AN@!c|QC4H=}z{i#E z0Bi!?Fd@e;dH!7LeGg+kzW0sqi1PgaUaNf9WP5ZWB9({Kdqo1yj*suNKj-EltbC8d zw@2(e1YkRD{9c@Rl!r>E4zBa?y7ApW>pgB);iJlT2sVapkWkN%l=GOIGv(v|sqr6B z{vX3b%6|(MMptyuzf$J?vtK`0|FbYwWfUs^Q}ERvxARg5TciJ8Bp&947xcII_}((U zOK81$xdHE2z7wzkbTN-_1@Bj$^YK0W^KO0P$~O~UqR_BJ;yz@kN49PZf->H^?ni_R^EBA9CUm1;VJL3($Z+}^Pw=`xdvFL^YM!D z7QNT|CHVStHjBe%(2Z+v>MiX){hKU#$Cb4E3G{nZLVTsa9p_>!%kRO*&3&W%_N^E` zL!T)B_OSfh6!Pw>X%a5KXQ+Y808O6htbx^WL25#?#5(!+k!#pyk&DkVxF9AP`)3-L&|px7DiV@xVb1Tt?Ye1*iJ68e(7MH&%h`?L-gKUJO|G{X|r0` zI&%k0gj~bqH}nyHM9)v)B3-f%KK{3j|Du|6xCx*74%>ec){m~s$G;u_MIZks-f-(5 zRsM(Jb;>^rR)8*>keqk6yYCh#5G{Xu_rKTp<|yAz_zo9VueKSs$UMRfA?IQ0y&D}| zmwbHRF}|~E-r+6ynDU*4#n3eo()Q#XFfxW&_VIn{m)$%xE8oZAHOe;!mWR$E%ylDJ z%86J6*Lir!_+}~J9{AR`+j)q>=9xQ~A|$_Umtv`RH}w1XzH5A^)m+2d@Db%Z4~wE} zAk6t)#=8nCdz_wvhm?0dEF0Y}a|b3D6><#?z;!NOF}{0hp5Z0< z`m=Ufao7yHaYAw-*ZlHunKR_$`@Zp=P`>ZN2bJ#Ox?!&`k7bSd_}(|ZBg*#! zctZKE$@b_%gq$Pt8$h`}@A&vWd&A8`Sot1@7bxEVY==35bt2W~%1VQk@t#0CZSDf= zrx&9LkEN`E(>-~oDT4TNBAw3p& z4L<)bZ8icMLDx@wIa=@Y5S8TtXNL7t3{J&I#>2~)$?}KrN#ifan=U+6kUUc{K{OEC zXc|=APE5Nwt5$WZfbaaUos$yS?)NhOA;K)@eYsutSB=*W>%IH=GJH~bC1F$OMu|a| z^*MCUcugv=yYMFEwFqlQS4YUc_jR@Q#`#?L@eTe*cYo`Z?@{5(w*Z!fZik%7HL$m% zHP+q1d>35T_loh|Rr$RPA5r6l1Z*1J7$NoLm9;#1>^L{zOXgneTljnYE9onlXv-s( zh)KQ-t0lU)z%3u<`6X!p*?b?&-^Elg{5YJO*%t1F`uP%3zTTC8Gj@YUkjLK*6t=;A z&j$K2Hr+utc*MO&nxe6_ur3F`QMOxQ{RrD6*k+!xlK8CicRspmH1a#PCh9j$ETZ4U zvxU0t@^=Flh9P{n@edFQ^wa3&-=^KcXO?|wB;;Mh8~6bls9+q{12KsKD}-Kx+!Cw?Y#UNd_!r^K#(%0tyjjrrQIw#Q$F6G z8gJ2iy+4NM{a?EtTd**?B7(zd>d(7CQuXIp%p9!ySs2ocLgjx7p7k&6GV5S#wB1EQ z`bYV#a8)?$wD|bmGQR1yPV-6C{|3DOU)%MbfDPaoBjk8qjGs$=mpbR;d-&Jf`o__E z^N|T}e8y&bY+sBnk1%7Wa^4Hl89mRrf?glT3&vaYUhlK;rGC54jj(0rJ?03SbZXoi z?4zlcu<~y#!M_;t_JOmp<*T{HCa^YV4%A5%w-%niMj3MRw5xiRY zZoq2Ll@RH9-pIK~U*bi1T}4D*SW!-KPWEpctp8aUt1=2zIj7*ONxS}aur>PmMMCPI z{6eIouRP?m`1sy3zDsIca05PFW#@YWHh?Zh(8W55bo(|F=hG1ynbO$Q-*uNKerdsvUeO7;n*g>w6i#R%x>YY#N=+Ylt`hRy+NCnG9K-5g+dljklf$xd*Sm zWanoE)_`xd@iz6AA&|4}<9%Yz&57u}^{#*ys(P2ec9|#GAf(&9F3Uhk{$(HHD`55dOJ4H|D#?>4FTq>uLl<1Kox_kDO*h0Qi$)#!?eRDaCBx4=II z;OzVOp872}A0^87S@=f6u5&ePlfHO`czFHhO*@Tnor@c;Z|16+S9l#hq_&)KU+`L4U z?_qd;pM4y&U^35JMdBEJqPPTC-WM{yE4?v7c+RT`glJ*@8(0#gFFcjEAKp54!(Qz#gb=v zra~?tC2*aOmyPefn#Z^V-w4|K5rfU5n;^`2n7TO|^zprCd?%IfUHFjlU4%8Gt0Nxf zS;q3@V#CL~;~wog3jZyt>w%Fnq&h>%l3C& zG{TmdmzX2c-obc#RNfmG_wjz)c#GcaeH-4Wyk}r>bj=>`Xm4wIfO;?cct5q^=0x;f z@5kXa|J-Idusn3~z&_J9?tN5cJar>g%sQQu%f?&uUhiJ`<}sUf!RDAVm?W|vdb{^% zgFfE(jJN2$-gn_OYFx4iYerXRyi?;6^Dey&AMc9aadRSiulLjNO*M8X6fL@K`eJF< zeO>2xTRL-XaGj4=jqi?{Yj_zxtbCKODRiTR)ke5xjrw?hWV}W1t@Hcvqzk9>npEnMf~HRHRf<}hA`52&%j zAZ#4n5bIgzo95&Av|=Q&luk#<@+3bL5&w`Ve8BjED?|L?DKDc^D5#v+8k~MrLHo4 z-9^aq8}Lmv*Dx;HPwcVU$1vm>!%F)hyoUUHbFdDpMHe6r{T}y{=2|LslxIJug5L; z#OG}`4NIVFdFWe~`aSNVkM9$2yY-DK-^0Sy7$*x>fG(SmcF|fEC`~f!TmaX3=rz7M z%C{4~Q*76_8Ma8jJ43j6aK{=6AK!P3@2u*l--0)&vBoSchOUW_W10F6nNiHLkMC2z z@8+Rd`92QcQ)7)BSROiuNY_`!8bxrOhnI|Rmh$a^Pbl9gY@W6^MYz|CKxH`I6P4cY zYw$O}QRm-vw11;6{mca0F8iD>@4;u)*kBp9{;}N_pIS2eue6WD_jLac%R?vs2vG&g z)#iFCV*z%_86pZ;UM6I+VwJz;d!SV!Nz}8A&eL@m(u{T0XJ3Oa9JSlaFl-dv0Fla5 zC=lZPMpfZnzW?y?{?vGj-h1wR3=ehKdDw!5(G?MLY+Fl1!PJDy9~`XnX_(ZxSoxlW z=cqc@z&7aHmregJzHI9>`S{*6zAI`ha2*~~@?n@_j;d!&#Ey(18jwMGfzmK<+pFrj=Ow(-!i@n zs^5MS-mHA5U`cdQ|=*u&G)Hffl^Uz~_ zbCqu!e7eNWO%rU1{&$uz*OHL*0Wyc1e#Vt=ZrVfqJ&0C*L*dQsE!e!d?#cIclSB+b z9p_mI%Q8&Q;_oT04byze_q7cykHFUm`Q61Zv4M}?kBa3~elqK*8=E@FQ=d=6b04!= zp=i;~|5Nwhm}u>5zt|n=ezE;xq%0EcXixAirby@cNC)$I^=M04?EScI%34=rhHLQ2 z&)Q`T!(^VLfwIO}*7d$`yhfDQ2k<84wI8x}&h!?BSnP#Z zYpc&V%J(sNK>0eb zLUilJ?tST{jvl@j>PUSjRKU9JsGaX#>M@27U?(#@a|&y!2TPr1vOrH z6F#K8reH~QA-v>zaxNLH3h=q$t)-6|{(yuh;DiwqL8{5^t?&ijWjR%MK- ze*7Lhq{bsFum*J1gyfxP--DqG^ann^&%EollN#k4gzx@NfK8wqCbFq7-}`j)j+G7;opB%UkBztJy~pB1crkVJvUON3 zx&R^f3GO=;%A8#v-;;mp<~XE$pMmEW*uEjyHphC6kaMx4GMIV^L_J*R;=1u&SL1_g z@MMwgI}95|H$Z#^z7M}cN&R~Sq11bRXMD;ywCdJ9smgeOw)t!AGIn4AbOnUTcL({F zF44)sN7uj2c#Gbf+tb38cRg&C{&s=z9Ltnb##{7W?>FJi-)OTbSQ1^7NIgH2 zdOt>4>Y}pfuJKmObFDdYHiZhmUn z&YPcO@a4a?&;26UKKS%a%puIh!#T1U1y!HTky}_{jOJaI|>iv z+jT2|WucRKja1!urmd2@nGjt2UNOGAXubP<89q{Ivjl7!-56o+?Lz#!GEw<`-H4C( zhsIm!)DNp6a6f&D6e;o*R=9_8(#PIw%0r?imrk9kmX@>J$*vz zN<)?3$jiUaJWLdE0JCU2g`D6s>`*q1Mc%p8MC`#w*rpy{2bcS+DfIiq7JrYjQ(>Zp zZ4+qwVR?L(u)gLm4z`7BFv-W3>YJ~?2cNX}VE{IPZkUj3Polj%+`;FVkMDisJED9) zfR`xWHQ64Wj1SmFr@Or>B%RcbkMFbh-1>%ws8$UqP&M;qv!?*$%VY@RPLE4eS9An-$CVjA0Ae|o3LthGCpv9 zL#37Y?)&(j`pZ-2v+#{mb}p)6o3yzJ#s`uMd8W~PP3ttVZld1S-!$GUs^7m3 zpZ!WZKVz^VbhUU-vaIv^!1c-)R9^StVO7sfST(w0!aZj4d)}T)7E0F8>H^irnuBPHD>v3<3Fdy z2DjnkxpqEgVDjF~X5v$KyUYQ#-N*2ds@E1QjBb}cK79?4Ym4*O2lw|hj5Qg>%J(EZ zN7bhWwm~1fOwb+h9b-bC*>{?Jd~X`x71d{7htGbwy`N*SA#@4Lmw*4!$M?{NY5&Ui z0opt@rrCi7&=nA=<5L!>tmthGwY4+;$44LgHsdXN@3r7GeD14lRu5aHk6s|s-j%(r zZBe|te7xT>-lF$$E=k0=z@8f^t=A%pbJ_fHpY?tN03en{l zU!EyTy~8dau5;02d~=m=8+`Yl*uG7$CHmx9B7HrO_7?Z?ecSlXsXqKRyitvxW?*r2 z%|trK6}>&-=!<+W}@onw24+BkQMS*P=H*?5cI>)i_< zSKeK)xv!%yBOcZ}-k$Q-<5GDS+L)Bbjh^_*wPKNw<-72CJkmTHCciUW61~x|2dVW> z{;j+3Ek>W_N8owi>1xs}8$G{^dy3Cu#+h||Ci%>Ut?=21Ciz)@XT^5g$ZJTz z02Ll!S-zW_W8DsHkZ8tpk##Y&Ta>Ny^UP){KbbYuP4bz}MG(IG9rk_{!}jQ-Wt@_O z*8NREMa7HmPoJDR)@%RQjsJ$~+poe$mH!ZI3|*u0zhL~QeEdH({sYSYV|YmUZ^6Ro zcImetJ^Nao`b|XE-yN*`85p~nS){z5gRilp-h9=<)}P>h4sQ+-&$(ARO+KDCjpqtl z@AlW>gUWLZHiWJgPYz#dT-Djmy^D|Up%2}5o>aaM(3U9Q9asQe0U`60kG>ZPm&+UA za`Do~vD5hHEB`a_{cL+*8el85@p&S3o`(Weo$>be&^g+nkM|wdJ5%)DoZN!PzunIB zG%SHGz*vT6?Xef-}w{?n@8ejDDR{O4g&bPYcKG8w+^;~(5| z+e*FiKPp`L7r?U6?a&un$MQV)!>rS}dDZpK6utLYz6{sjjU-`H=%Qxa)7|U-TW=#i zz8@OjG39#?UjJ=&&R1X!=yLFV)J963Eg#=!{y(>k)F|H|e7Dy2Er#vU7tb<&NVSpn za{cx1(C6H-kK+z3fUbae#@#Nn zMr{{_@1C%?D~9bc=`)*S^&hFArQ+ z$Nz!rpE0QX@58&k#4d9aR*kNh_$)blm{XRW?Spkc4P#BF=)HA6319zKJ1;e`4f^cy zFZ1T5A|%~*vybF#@gU@Ce;1DubIzeo~nI=ws&zjAuQPwrGtfZ}5M(IjL8^N8wwC?6L}A zS?G35-z@LF<5)T&xX#5Drn}=HEdknsz`uif-KJyz}gp7SI1xxuh&T-!(ArIA`t1Ojg+2mL@ zpxxyfAjgc$L+amUm*rRBS*omln9k$7FytB2YCh#{_{hF}KFq_S=o$#QM>==$;<*?D z_mANNTnD9GIZuj+X{ikLNy8nhn(b( z_haz-|6_YQutIb>gxn9h^{$fd)g#H6;}o)9AEQggKU?`wO14`oOOD9IOjnBa!>aH&yS>i?JxkS@$U;_`lrT)T=U%!UL*|0$3Kh z9n+_Ja^r+qr}Ofv@fN+e{x8E<|F>QLBy0-ZDDmi+hEvCM)W`cH<1Kox_xtdq@?M42 zqYD#{avpZJeSA-R;^w1T`BuOSm2U}b_p|)|hmbLWn~!kn84PLj51(Uaxe4tM=bY44 zPQ5NdmS2VMv%Pm824QpRntl(4JVQStdRKh+>|8Fwn(?e7?z22-u8VT58sx;{YphiH z3;ZqR?xNAh?&ROQ<%Lw;o`KKbvC9j=wwd=>Bjnie%u-)fz}dCQ!FCDBHScUi7h zW1}Tl6YFXTPx)c z{xB37g=(B}3ckwyknD?C$Y&i)o&jAXuyf+)_#TW`L$vt#-ZH*RYMgNcKD}%EPQc_D z&=?`r=c)hFvfk(0E6(v=G9T|F|HrMf=)Ltj1aDB@`)pr?E;sGX3tPjz@z!>{^U>Ft zeM|3SkMS10*SihA`#3US2ij?5O#Lm*K1b!_G?*Hid4K`04}Y#O38g zsRslne9HO6lryBt`3N3U`?>T;3~LGNPaJM@SQ}c%4zcPy=i<`(0b4F z>+sqCZs%bPHiRxg*!Sb*skwwXAOFM7Kcwm%SN@ss998cJJ_6tW#AeyBe9;j)-|duD%zB-d%f{bP{=M*V<=+LHV@xnf zq_3}Xz6|*IzGr-=)HviFc&+kXfVH5jC(`G4IsMFBieUT9J!}zDXZcNqybHS; zuG_>_A{gI3WPrVGd!Pc;p>^fnuC`*6Dm)H{PQ6dS8PlmG>}g z6y1RBof^YW`FMY7yhZQz{uo}MytiOsbVWq^JPC1~cE0prolnEqos44TdlJ5=#tSvD z4dxA&3CRcTjsa>}s@*mFc)w}9MenWi8}KgWJq{a07bhOoxzc$TK4-QUk12xg&S@m% z?^V_(f5Y~F`jER%jjG%y;XA5t&x7Tl+vEPfkmcqf+v}3?%2r+%;FHR$1-5|CG!bKY zR(ahqUh}N?=HnK;NqJ4f66jipIhF@i|8V#-?motqS0+41dF`=%F}gfLKYvuo&nEKO zril(r>clNo>UnQjehyx3{JGsqeP7qYryPTN*46Rp-H+=qP{z9If3Lw8{)??0hK-^d zFxrdp_(isLCgEAXZfhSH-$CVjA6}@&KAW&=bj8FD%gxH`*=)CtVdZrkzN*G00oV?C zUnkD6oNuo0Uqe0fvpFs_)EJ+;4Bi@9uI6)&zlZr8<>D#N-Y<~2F#0S$ceuIQV0oUT zwGaa=%l&T+>ymuVv5h>ZFZa9!Y&%1lI>(1T>()Q1+QS31%^%u%+kpkp6%aDt`TY5w zbI)JwITuN$rm?tpKIq?R{PUIn8TkI++5Qc%75eCTf+{mo>n6Tfh&RBrm_eSD8R=H?=%d=J4Jlp#nEt-GO6u<8{kHKIJxzd7ru%stKCZlHU~zQKgqw@5R4$f$ ze4qR>Hy16+_XvDnjUTdMskV%krkZ-bRPH+p;W`(WjBmE`y#U|-TYDc`U<>s7GOr=y zO3(9F6{Y+=;N$TG(^K_h(i?OV1+0bxS%VAIK&|eQK>@}qEeNbKf0icbOj7tVMd`YRj5l`s#1q4 zL?sSUi9=kfvJP>GN>t(yg{ah}Ds`w!RpKQo@v>gl%X(QYtM2!?=e^=R4`()0>GOWh z;rl-C?|shmzUO`4bMCq4`sjab^oJ;8$K)Yge;;WHR*7AhBr)Nc-o!iZHhgqn{E{1! zQl(o0U;fxzS3YcuIeL}E(3W_wR&hy@Qx4bhxM_6P)SP}D-m7%`U?bS|dvpt&5g*-8 zjqZTbeFV=_x+}19?1Dx&aSw{Omu>mzo<88lA*6I)fM@-gHx41#2J`7MN#Y>7B|P(3 z3Dy0?w)qMDy?!DC8y z7}k$n+~~^tPJ-x8`snV|0b6FRFh`Q{f)C0=1%(AAE$l~pwC@;gvA5gqHoQ`4kHLDdiz@94XwUj+ANqU^(mV11^4qG$h zl&q^qNpfuc@IJFIyv*!>1ZOueqV7+80MAw9(gLg+y9|6g{8hIPL{zy0!pXPenF`Cs zE=>}ZU6$`RUpmpju9HKa?gL#$TkLI}i}1l!FRO>mveuX&S=!Qd_gr_B@am{R>~jc! zt@|Eq@-JNWdCGN@7Dj zNniA6|DdzxqkZzb+}Mb{t$hr>`jMCA!glDh>m)h1ys_akf!3xruv3Yhj?XPq&x+~u zG*zO$8}P`Ox1Ism5O%UY*jrB?jZpZ!)bk~hJ$@AN_dI3etFE4$Km8;*t|!9(o41}V znD+DZp+r0aA?4=_@R|SOEgyny&>zP+{|-~f<)%nWL%iuy-Q{?=fJY9`M?2asH#N4k zU2>dTMt8;Z->1h9-&ytQ4#0-6t3r2}{qA&c{(acE8TYDe0-41D#efjP3 zyszzz*{d+BNO9D$b-!L9CK&hved>xZ_jQ#@U?3la-ul`S7HUjIz zuFU8^9e3q@IOxy#=pX!s>wiq??}KkGd-ZoHpNrk4H;+3 z%>L1AH!jtblRSpG#P(Tu_8l+Fgr#FQ|DX9jid!xvTa`NtANx~pxk}g)^YSc7`We@S z1t<8ndIYXxdB^C^VQc$y8$R%RUN#2n!7gg)hL0b2W_)xHes|(HQ({VYAAIWrukH@z zbFs@HF%>wiZI^hfT=rSz(DpTg6Yyjtt9 zFm|~liB;m`}7ig<9EEQ9Ja!IJN&0S+vb)_+2C)x+-vaofAF#% z*cf($q^Es~uK^Ch{o?eAs~r$~Tl*t;jq+(3Rv~uy^xXHlV?vE8cNo6Vjvh^Ew8b!mb^D$1RsKr^?-i4@`Nr#$Y|zMM>7$pq2J< zX36vO>+lh`{US>L06bsmr^2$aTmFxZbH*)~lCR30f=~T>uV1CGHRk3;lJs>me-_F0 zNDX;<%)f24#omt9E%@+nd)Y9oAGv@!4=OD%J0BOwZzYNc0d>1a)LbIQDzti8okI!Lg zEJE#Mvz_7ZQH(MvET56p!)n>C!M4d;!>%-iIh}1D8l%{)W3$Ot-bp3jRgm}n2Kjr0 zzo)6U0$b42x%@z`+ve@83v3>MSN|I?+ors%AJ#Z06YKky&ZgGRrn)A^^i1rv_C=$e zskC2)Z{PQ7SHl)qJM^P1el(OEmrFG#=A-+b(VfBG_W51-sM4K)#jy(-pCg@h=VKvv zbe;Fne)jv_c+@NHgYa^tod(OoZi(?g;&Gv)p{phSEq>xP5p?C;Yi)}+)fE&wgFd>CjILhGdlSHhNFeu(xA%8$N!| z%f?_m*oDy*+w+YjVK&Z;kM6x(%qqaE_SQ$yPwrxEDSjr*y$KF8_g`G zc@94D3tnHVV2jM}188#V$SLGE3_=~PMTIn8%t!k@qdh|zI}Y!{tCaQxERJ2!(hhgD zmYl!9wVIFivq#-n)GO_S@U3g!`qE%o*e(6rUB46BEvy;I(=oYfv{RLKD|~&rq_neOY1qvzxP9Y%$#EwirsHtUXzr-BMHhVa5530_gH2)Ap42QZa(aCaoj4Q^x0p&XVFQY?pm>kL8==SGq^wA*GuIlWT}M)(vtlpO>Es=LOzD zxQ@YVMt9rH_sO~HDtzTVZw%UDli0=3{havAIoIc-`JvGqQ=0F?`<3P#tQNa0G^Nk9 z5(75Qvp;0|ztT7e-_i3wEDO6OcqV1l=UPqSqQXMa3}PqWosoK2ia7&jzZss+m@QnA zeC{>JHjPwEijuG6b{9rgV3v7X$|U3V4t(qnylfmM>xYo>J6;fcg(A)@yo#+|{?IX3 zH>z~^!?TsG1Ixi~g?V3&r-;bOIV$s6!wJqU^0nVxMn6sIUxaVWdVQ&f%`zAFpfBx> z-%i{|^S;s4*BE!DbzA8;v~d*533%x1j#GA+e)T!B)J~MV^E?abR$*ZA`dX8!NN3p?=75tflPaUOR{vINKD}+y!+n^uTvW=tD2Mc{4 zeiu@{fI3NI#BG2WJm?}>a7yO#WuO4)l@L`cULZl5e;~e`8 ze#WS)g5ypjp6U3yf<~Aa)v`ZD?6QgZIQ5LvKSIxJ#1QD=l>~erP%3qkG5b&S5WpChNNmUz_u?F<1|FVRQ+T z6X*HdX&=o4MdtihngO_8J8V-v54&VtvL->=$-q{}pxM<;$x^!K;0u4~^|cDNxImkc z#McJyCq#TS@3@))SzmaLAHJqE$6!6!g^5cVntWb%@p#ekxcEEkqkSmk#v!V-_rptn z!&{#N%fW7iI@8hHeLgWrPiFilU&o`%=%*?Di|{<5ogIq`)IKn6MY}#S$O82QadgeUoG8rj2Qzy+8-J15z5#x z`2ari2VS-StHv$^Z8`rUEf-qqItraNAKjBBZanmTkYn(O(#?hK{0UC8KEXEcUwd`8VENc(k))r<_YCTq zi=A{dD)`&>tI_CYE8Vm3#b5KGG4m@$d;3_9-f2$xIbe2@BQ-*vSEVsF>^4t({l zmyN@^u?ul5UqJ7CfqOkN>!W+l zPFLE^@O-6R3!A&=IKA*R^pfWdyH2-{_WMSAk}`HY-h=n4SWm;E*yW-veXpabt1DJl zn0)lkS@zLCcG8WDz7KK)o~!h;U}@;jaX!g8Dc_T*OMFp+_aBj`1pgip_SwHIMk*rlVLiQmbw#GTi;q1%*r zjMmxqspq91apMzI^}Gn*zUi$e43qmDW7IEgOTTF-VFF^;sU}ay=WV0CL>W6iZ^B2E z_7H3Uy9%^TpA_wet7uR9Xzwd^{q9!UkFl-%Rd0Qpupo9D|H{>NfA8#SA}HzTXulhc zcDB+!3*Q{^YFEORn2!h1-iNR5oXkArqkG5b&Z+gnZTQ+9ukIL3)(>HHKaXy4K`8N} zoEab8gFouVBBpfr!3UJ?4&`P2unOOgu31-gwbP1O*y~ud8tn|F-3SjV?HZV@ANt_Z zKN>nNw{_HYC)QMPAMN`_drGYr?!uRT&l`^^n5-Z2&`yq7tOpZg)}oL0b1%3tsZrX8 z;W@)z?F^W#9~P)n{5JhMvDVBdPsik%(cV#WdKY}|m%Z9C*c5i{Xj}a{&ZUsk@1yxs5j-um3fHf@Nx z_jpo&e6QcUa^-h9Ox6#>@Wk5WGQWJuH;!t3bl)+$3)tK7cnhwds~Cax zVONH3V%!Uc!r#KFDIeYaFS`E5mF^R4hknOf*A`6HD(mp1ZVS5UXz2E9bahj*mF`*i zqB>tHVN1-rvm|N1#JZ9TU>Q83KH7I(?SRZ4oolWr_3lw=$+AsZ-8;hXw_eFS>(hkGc znP11?$uU%PtKd2gZyDWX?Cm(*gm0^UJO~@at`c3DOFE-GDH1&4jQi+5F}l4<_fz<& z(p`szvD<=+Zsf|Px){1n#a@5U!Z1n6Q@W?&E4RJJPzGCNzFr_npGc1R>^ikR+V2=` zvA1LK7Q9z!kHGq{YbPb{UYh}oR9Y5({4N>O8Wr3Pid#Zvawr+@5UqXm1!r3JROfNqn)O-FT(Sbc0FvC zIlBig@wgmqNc6jIAKmwj?xb3;yax~emNy>Luqbx9Mpx#j66UByAKmBviW>|4jObzb z>bO@o1D1jA!h{Kr$Y_77mHmDn z?T1EtRB1ne&;E07ee@s)#md^;<#qTvA?UR*mOv;t^F?jmxUhQ1i4)f|XbxIq$ z>w$|-DO|_n4Wqk(y&aF&L|6IT0~^Dx9-bUSub@BVqx-SZ9a6dv;ZtvW>so?UVwZ+) zn)vH}Ued*`v+1LK>aV)72rKQA@EWBZgl#gnPQZ=!1wMI@CUh#v)A6`vv{xu&$KwWk zNA>9e*bsJAXq*1gdA#I!R}AfOAMGbbyH{y{3LjP4>##6(TkvE|WL+SBr@pk;?{hG+ z0=Y{2C3v~gE{Cl!#}4mm6BG9yd%chLyGDCntpVPKuaA4O2Yd%);+ ziNAaX8uih=YjkI^x9fWczWSS9HV*5?E@X5Q<0!hbKDvjh+&Dy)?tXZ`(sf`t*sZ`N z4xRDlM0~U0Iu5NyH$&+*!iWEv*WVi0JZp_Ube|Dj#!;?)F35=z_tCy@w5Qa%ETf5o?!NSD(oll;Q$u*kHQy!-FqB4VvF6(Z|uf{`LI=DT}qyg#~Vg_L#+c|gU6M24{Qv(dbE>$ z1np)KANJAy#ApvH?T_GDN_$z#i=CxS|4YoZTRz&Sf7XpjNNK+SpBeSWBm~=Fo*f_C z^&9Pk-&N%G;lG^E0aBUSf7|FUQO-Ws--M4`^Rgk>0CpAVf8>@+8Bygvfai{S%PqjF zu`454@l2d^>*V?Q^`q0%6j8Vpp*Zc4})%WLMwb*6h*TJ7l98XG2mD>m3def`5L%Ce+ zGDyjI(T-ON@GA>@KfgL${eWJ}G{cvAz52DVIp*?S{C(Rkm$IbR6K}$6m0v@!0qnX- zi81+dU3($7DktFjd<&d)+pkCIeuizT(%pcCu-oKZF~{4<*cOR#ylLoY?U#*qj?z8@ zAOBVFu~xvAe~Y!-t~Q^u7bapE@zK6xwC5;ex9@GZ{>*9&)`MLb?KFIqxrEQ~FR<&( z`e+~e>uyYMNheo`|QVW|Kr~N(JiCc?K zw~zMwMthPnc1+%bhn4m;EQ+0s0crSZVse>XXURwVXpI|_YNh=weC3~d>&t|tV>i!v z{LHR)VtfmdR{^vAd(CKXGauXhD!f-|x5MPQu^8H>Z*vc@*y)G&8QlQeGPD1Y(H>FS zAHXAj+gslPtQxxvv=9FccN~bQatGiWH@)RjVcFQFNupk}R!NLcIpq2Ib=B1mq$>SZ z_(G3YKLVR!Uhd|6e8(-9BI^q~R&T+pm0u&UKI~+@kZcFW&1PxjY4UVC960B;vDn+% z0r>W>c-c1P^RQdvJWd|7oUcebnb`XsbF0zMQ2LGV<)8QJ*TCkPd;8FT%Pp6(%-^aeg{G5wCLc{QP>;)eo#vCKdlvNV0tcK7Q59`eDP^RpZwOZn+d$ zU)bgD!-LANSy(-Gm84`JVGVU5PCYB+^_e>DbE>4)jhC!1Qpr#HbrhbV{K^qq>}EKR zTitRg8LC_(eE65WS~akF=Hn^S=cz;P8Ju9-4)62v>AtHS5PRFFyYQZGdD#>!hFu;$ z1^!mzSW-PFp30q>WALKl~ z>6S}b<8S+TuES$DysQs4f?Yo;8LJXvGy?bY=`&Y5AojNQ$M7oU(<-bCyB*H6V=ue0 zs#4{Sz?c7tS1SvahTYcgoGasWV%;Ai&(E*xu6{t)7q(y5;2G*X?}m+I7lGe%%cV@K za(CdpT-DiHldyK|YDuOIFBF{!C+4OF^8EaIuFmc6HA??5JV)tgz%sB~;5>WTEtis` z%AJ9a4|)BnfGv;mStRKTv>kH@Yqnyi9vq8>nS^Z?gH2&KN-}*!?j0x2iGCmLheliM zZJh`3>RJVlZS7<5ZC1WE%Z2SQ*RGRx{Vq;umyxGq za>Hngy{+>)eB=!;>xGSCH?XTM=S@O;guFg;4%_GRXRdyrU+I4gA6NRTurlm+;K!QW zW3E!=j=;0N>GdlMmWJIHbGbQYdA=gySBN})%&!}5vA2);Rd}ZIDGr;!ZkS~HuG|kv z%wdB*+K-I3*xTA4!l(bSmo37ouq(B+S*IqnH^}pg=c(pITmoUGe-d7!^nE+uG;g6W{c*D%c`xff-V=Z;Q4J`Z4nQe9q^4u6|%zts(Bfm#=&ECt>Z_ z<^LnT+vb)_=~d-Eg;&smwy*23Fm}15WdD;_25cyYcauHQz;({1so(H47K=MjA3A9~petQ@-_ zDS6DIEyuh~o?ly>x#%8qzS2JhuT=V_Fj=dN!e4jGrEIA+z^m{XuHJ25(H!vA1vhw!@c8Rq?MzrY zy7TZxw_HlL^64ymE$S^-36p0)XGzKH_#%9Zz;*lHF}ibVZEzc2!*ST{G6w6xE=o%F zX_-5vPtTI)cf5yN-Q$fY{R8m6U-asy!sHpyW%wDlTuQzwcM8768EX4g3R`2oT_jok zj^4?&8F~7cZyRl~w{>p8H~YP87}k$nd{X8}OaK?_~q9A?$iB zZSH+0j#=tT{=Tx*HI403id$aR#_}$3(fEYxDJmID4X@cMzE_vBa^aw-aR5OUFP2ec5B~o%SxMv z54-yx!3WK_`s47quXEkEzGHLrrbewtEHsyO^@-FQ$*b4R`^fJ+s_Y=r(42R(38GOb;r)a0nIa1aE ze5A}2Hkk+coS(n*u*sp`DfII`M>)z)<5vx7nL4^jJ+Ni$H?XP4W}d%?IF@C~=fLvW z-$Jt@i#(1u3)?(2alz>kAMs(iV{f}2>va2hl{r6O=34f;!qqu$Smq?qn&(o9~#|JrTYM0 z+UM1shsj!{oRpY*+Ai=`YTDgLwtdF7s0om)Lbm z$&f_k1dl82Whswt zh$L}leNZgFfVk z3T15V8>0O!FB^ajVb@ENHtgcJ1Ve1cd~_ch-9DxJ2|TKF*I=dC<&jJqHlHXgX+@e! z_WFDlMox;@+diL$rz`C;*y=BGzm#OOX~Tk7M7q{T`yHb#_O|w0@an(qWg{?ItF)7( z4KK&##}yene02A};`$s{x=*mJ>GtYw!Sb=oB3bdcB0guJqua3AXp6n=^EtS_pIQZ5 z?C1KI^bC4#pANnvlS0f#`#qyA_O|xB@S0vPn}Ef!t0$#B)hmn;uvzBg9_o`ej5X1Q45ZyY zy6+p^Ni~1J2M;UVX;>7yTy&Yj8-$(nQ(H47K=K;J=Y0txC z{ZLNY?RNzUzt?=UPhNFnBKEfSF?gQR&V}v#J?{UKq;Iqo2Bis|Qn-%C8%B2nTU+-v z_+s44dSGMN4eshDzJxL4qx-SZ9a6dv;oVAi308?+n3Rn3@q&WHd9vxFed_1kSctvt z^GSFv0yYc6Hv3oukyzq5omV=V8_qY?MXt2TD_oo^@^wt!a`gktl(DsL!t2%8Gzc5S zu9Ef0KJ*gjOEGIJAKfQLw^!+Y3QtqI>##6(TbxJY_tW3S2nNFpsZLtYUcX<4k(H98 z^v}RkJH5wJ0bBkB)_tVi{ujI=F^TwS-!a-^pDLA*WP2Mv-r;3qupaE9B-5VEF~LN8 z&iZH{`rEGGVsC5jho`D>#DV2tmrmMG{fXZcZwc~L0MCDsr`xi_)eeZgt=$YCQe#Oi zY>st8FKd%zTk@JMC+?$r-{?+ZZ|mNL>oIBy7Q-&j(k(!D!AJMVtBF_yYL)IG_?kM; z(_xv|Ei%U^{grQ|=fZUyt{UA`rP~T0P`VM=%->~iyISSV?yzApSpTb^eaSy0MFp?0y=GaL)N? zAMSOJRqSo;1HzSdDl8kjG?L7v{H6^*-&g2l!?n*HMmJsQHp9nR_1SgR!seKtC(Rs< zZcz~nxVVq*eWN>tt*vnvUZr%WU@`1!Ji3L>f{*Tzzw5@KR_PvsZ>#e)9hQlmLrP;K z^E2~CBfCx>c{&!a8f~$+eeQsdDD5b0nrn$M(ihR&9cRuveLmVB8f~$+wcm%ADeXB} zEp`iq42jbS%P+Vwk_i1)CM_9sSL>}~Cj;Q2~>S;}J<+SQh}OkDG9k;f%%LPN%fFdA z2v7Qxtm~1{9ag#@!gG}FBCHC#bfcSCKQO_s!~2Zy;%}MRf3eR!=2E3y0?)bR^*J9V z>xWg+RkvIqRh4Un4{$%jKK=-7hBDo(8Itn~KWF3g!1b{{FuD_Jp1ubUDcu=Z1iKt` z(>b2bXhX-Px=WoMbx~#)S)=(MbIIQ`{a@8_6uzd`9ywx*-3D`Qt6MIRp~^MF_2;)W zuzAY#p(pb|C--BD*|x*=G2e4_Q>L-E<8}vLt^ApUwPROHN)z40Jrj3ryFi{_8$b7& zd%QJD|1f-&b%R}J227p--C<6B*)12yQRU9S$1i)!Rlt@hGxSZjjk&K?&>7`g6|Rr> zT~{|{9(%j)x8V`W+iVopi(QN)ag=M16Ktn_bPxP}_c+^?ZUCOHbhjy=hg~K~<~;cg zV?G;oGO*SDHXGe6rF#xO%6V`5TLoKWew`sn`|u%OLBXZM3+y^EAMN*yw%FS`@50NK z_5>`Bom?}>SS;FwmkJ6*e4e~MwxVI5-$(irF$w7P#3A^aI=|ClnP@L^o}F{crQ|BV zUV=|>?%KYV!&aE1=SkkVt$}+L{(jwY^#gNiZodt$RDO-Ydaw)Q*Qai|lmS)l5j<1* zwE`>0PM&51@Gu8=w9~x9^0M?hEimW)1szL$D3z;AN7%7AuKx{YReGe#>Z! zy{&T-Ue8Envq9J(c0D8+)8uz=>Mro@pE~#R@kt;3eXqN|cPstJQm@kAgvraYvq@hw z`WI-Et{QYMy=* zUi-t|azik=hUg~Ae9;+o@7+$ob!-Fwz&-XJrTZDSS!#^hfXOvPPExnc)lEgi?^s`U zwF6>rx6v8+c+|@(V9U(8LouHJaLc8v^S52@b$B27HtU7SHN=3KuRHOnsF7zQ;QClU zF}j0F_ak_g@@HAfV;3UHSiw&kv-~OQl(FLA!0tU08B@zyXl${qA*GQV$KN_PGoqcumq&2Pgif5^+mU_IDHNt5i4D8Kgo zL(~72)?=xI1Ge=yVL|M&NjvP9sd?dy(aKj^r{Gh~UaeA?tQQtZ73?n)N4u^!jn*pp zHh%*i&G)i?*f4goUidOW_Wb@-NzfT5f8Z-_85!R*OgtaLCyc(FD_Nvzl7la!q-u66 z*{);Pt^9c5x*NAL#S7sp>Rir)Z4;|CQiy%sH(xheTPmKf!h3$ud#rJoJU2FMY&+Xp znOQjp;Av_-`;pNdQMw<%!^+16ST%OirpxSWACLVb_js$6))DyfQLm3#ur%zp@JHqj zes(PJtL8ztytCbI&)1CZw(|WdyjSVA!zQsCA@#67tLBymMr%T8-Ghge)(k9yT@~pO z`vYdYdmewsNTZ*lKjD0klRCs8Aio0;p_8nUx5J4)LXvmo@1!%gvY*9dmPwf&Ch-yM zrVcQtv!6--3A4Y-HcslH%pz$E+X4FF5L@}IxRjvCZ=DUW%|&aO)J@8u?lSUg$P-;z zpj5+p2KTnn0~i)56U_CbY;~nWC&6)4TBDS}&V|#jvX( z$+;PCEaJ)R^Ui{g?vY<|{jXKJhv3^Zq+M4!EE7A2B;B;N`*={U$#dblt*;v0RHfSr zU;mq4-3V+Z#yxxzry18U^dh!BKDrN#?gX~BZ};GXN_Pe(UtX&sahgg04;K^`6kg#B zSoYCAHe~w0(mn#;IPa}53zml67HjdRv0zRN!gVZOGrHT@+CE=}4=LSt*d%r%PjyR< zON{${bRQbsQKkC;-gnMh*F3BSyK<8B3Axuu>p81Fx+VX_jYEafJqlmU_3Gw`Ep{9D z_$74BJ%%{H`Dtl=xq6DpT6P7MC@&!Uw|(*cv%Rx(Lx_1NlbRN3+fnp zoNDrQOx`y7OO&zo--Nd-{UO)@cJj=p)~7o;Q$G6pf7$gvuJoT^o1yf#VENc(ktF6Z zu6$a%PdMplXulhcZnn}r3!kp{9!n)`Da!aqlJi{d`J7-I@zK3wbmy?O<8d2auXM*? zJ=jH!Z9~xgeXAKC-Gjq!9AZj$A3ROz?od7#y9|?>!Q)kRJt$2 zN50oQbkmh?GkjdFA!=cBoDh>F6NivI>lOFWy>E2& zHO5_dmC~Jp#n7xVx^m4}KpYl)bdTJ0<4}vO9p6Ln?Z4$^>99=f9FoL==a1y)jGSD! zj>A=>o2qnM;iF170-I@K{v&+>JsIPQxdGhcqx-<MdrxeY?hGt~T@@(}UG9Z; zo(Q`4LYIBCkNs0O7GiH}AAxVEHAWUJ4ZE#~t?hoM5+YB><8`Ag_O|w`@Ij>=hfQEN z?A4|=?*IE}KQh{4Z)<-D4=L?MSQU1qq|c+rS|jnS#=4L03nOkk%9L&)d_}D_@?hKa zu{DxsZI`@nWY%_7*y-53W$Ia`tnL3zcwE&p2phz%ha`O^@tGm3rEwqKCq}nd>3#~& zR=Vr3Fm}15edx+L6Jp=_)xG|ng^`nzr*u!lC$qiBPzGCN?p+`y`*w&A<7<3$ZyVi3 zHCNw)$JBae7}k$n+|uRwcfL!z*WZ0_y8d=6-N)EwD&0+(e8QGZk~p+B(Lq>@rlFzR z?Pa5zqjb-}Xa0sa4i&Iv=Gi&YQ{5uQY9HNqjqbdfm*0k0eZ#9e3hTu#M*1H7^|?+e zzEF4Nl{#~^wCGdcbHC=sr$*Ix7`}SgTVDoD)(boJT=(v&kNczUPfYE(N5l8#ol44% zec9CuP`2LeUxRO&vFj*2pG8f3I`bn7vk;wTmi-~PTvPWPaGV^{0%_-S^j8M9Ic%b& zl@zX{If3S}&3uNjg8dn`t87D@>|x4>2>vqrqMc6pIQw${rwWIv`FoV@27k+afE9H5 z_ER4ka=jjvy4as*Kb8IL(Y<}514gkxdc8Mp&G5~aysQ>BSLZmrf8D*8-B7?QYukAN zY}0vWL5|f&|AEn;pqxEc+=CA){TWyUyBzd!nZEoAKOyb5(6W#Av0rz6uTt7a;31`* z1xv$j?yP%WcZn=dcm?4V)Mfkjn$g`>^Ym5t%3t-`w!fFR&f-(S2leht>T5A-qTFF2bs? zOMj|c5ajllv+kq&!asB4P^NSX;kimT54O!*JH-vvguk>~1l=;Yj>8S3yRPQu*Wn9) z#T$oS*eG_jN!^ly;3ag2eRMxDx_UkF5j?JRm!-VakFNC93&jOtHqNGx?x`_14q@f* zNqDx>4Z=2=TPJ?jJ?C2Z{qVZ>lH)}&#u@T-EN;2lDSAC|13r7&8;b$h5Zb*YhC8Q2 z-ksfD=#2U3K6Z5jec0OP;wSK^(p`g!_{EnxPu5p`6>u-(Egq?y17dCC3u?BEr+eJ9vDvQhKl5|G^f@__Z_3Vpw z!6$LTZp#r^A9iKv%2=5gdx-CpkM928bmI_Lx=*l`?=soCTd;iW)~n2Yv!a5+#`g2Y z1>IZ=qoZTdY_zkK_Br^@PkFVgV2d@}&p})I@P%gS!$nTiNB6GLoyFev?;ZG<(jA9& zV;4de~C0*$@!#7WP zZEIn3%*~S|iNz~zkwh%wKDzgf?i99meRtu*N_Pqt!>%T&8*yX2;G=uww-Rwksa3j% z;AKiT9hQlmV|3l|r@+aD>o{CBx~WRH6~6wH-eZWsW|)`9lWU3ArbHZid~_ce-3e^% zINXB|D%}}a1iLDu>-t;dEcxgjecO#gwbFeS9#*=UuypLUnL~Hmt-#5L>o{C9x;xm~ z{&vAv%e^cHo5F51sheoGLZ{D1_d}yQrgYzj_bc5wSS@xHMz<>w-vVdFN4M~wyK$&g zy3fIbN;eypirpr2=x!Veoe*5d;dNIxuw~}k z;n~V|6;_7b4s*9$lNFUH|A~3(zu4>lIT+?CxvGwr;At;<{V9j7Fjvo$q@QpfD$Zvy zPOXpbJFaeE0b4sRZ^0*j+{;E_eb|-ZyXYpKQ41G4Q$D);Z@X<7SGrHItx>vLuzc*+ zna7j5A)b9pM?<$`qtVS)x@Uze-AdRJbM>I5Tf}D>5g*+4DMz7NN6rQctDC@8=c3T`^Cb8V}9cw;91Uni3k{G1B^-F(QqO*&Cr{MYK z8u9D=y+&q~^SYO96*klC%ZJsI7^iaL$!G1=Y{yAqvNp-9WF1lsUt^n=#(IgA_Y7r8 z`COaNk|NkHP$mPN5%Gn;rGAl@?WZvLQ`Q$Ye*|y;F)v$zm17qqA(^-rSj2YANB8uk zt1IgZTlWR{;!k;52)4mmLDmbR%YEX+52RGWwZFHF?h3ZH?hSai(j9;eVb@C{OzwSR z-VHqFqx;zC_9@*@;5kZn4OWU>9*L04=cJ6dCoZtdZ$s_%`7Df_6tTB`J`JB!Ylbq| zD*a-CM6)`)q@u8`ow3YE_qNepRQDQg!R6TqyRKnaKX!4F^jGhFFS*%O@I2U=^{MC3 zJFfpxRnLBSs`B4~~|UcG^KwLKK7H|p7 zyQ$k|1F#|NdM)kaMeg@Q$9=S)7;Uk)bv}h>EA4ey7(2PI;hwXF?vENx?e+OAjKfLE zQ@W?&b85X&23uv`miroCppg67PO($tqkG%vE@Er@dkY>@y2G%3?Bu@2Q(eaVNgv&P zzvKGbt#lt_o2hg+VL|NVzJ|m>o+p>_ISmcnb}t*<9Hn~(KJ%jY7%E`P%+qruSGS;$ zrG``Qqx-JWoyXRW!`tvGfXzl>z1YQ!u8g$>=uX4?*ov)vZXSHsJqN_z*4YQ&`cW_2 zK_eHt)gO19({8zxkSg~Ae1>_+)(XKkewy*0ltKO8vB{lBYp~Jrylv`Oq@1mH3tpq@ z7>4y@7bhLwQ^)Qte;EbmeCj&z~+C7&lgye zyycclS>|uM+)a4B@@Wt@h+Plq3&bSx`JX)h3itCV@Oy557kgX#Gwf59PaCihcAGCa z&PlgiO1Ua`48CC6Be~|vh3znJu9J3sDsoEUgu+wr8%B3S&9kq;e<&Zde>;plW(`#+wh3e8in;@7c;h9EnRK+ z;Y`E()YyLDU%5WEE8PHmn?c^LYn$?U*ky_>`P#>GMk`loy#$|7TIH}6;x$j&hi~02 z4Q-cO+_-w4;|>+Z#P>M%tLhxSZ|a*u!}jwod|0t5SPWJ}nq_}eX&stz;}BKl_QSL3 zEUEaA?As132fK7hQclIc(P-)Q#984=zY?}Yo6M4G*dN{R)+K##gcC>Vy$)Y1bN9cC zzpEuv#(pwqKkqm>*j2MXgo9gbYYC`4t6WZ6!n7aT9P--PpJSWJHj2g|b{XWwD6>Qz zx;>Bnej+A;YJ9`w>G3}cPy2?KWx~?2n?LEEKUbK@^I%%@Rio)B%?@}DW$k*Ru;~|R zd&A?$6OVv8y*|1h7~OH~ZH@cz`4e6?3#-R28(j|H;Rk)3B_GYBcg^uD&1d0prI`s! z$4=^!{@fXMe>68Au48b`=2B_CZ^8=$ft~Bq%=heA4 z3#-RY>I$G4Cx6LD^XR`e$FDS>g-aPb&Y7@u?2>gA%a`H#mk-}_{6=@D%&XA_&sMrI z*c5i{hF>}Ee%HCrNB2XcJEnBshtD1N)-?yK#V*U}9uF1~4X#mqbPH$QI8-X#=io7= zn+;3FZWhj|>~uA>UiwyBopdoLOrDO#b)&t>oNveT8azX3cf-c9i@>GbI`MtLNAn}2 zIifT_fVXFP>sf$RV<&Z`q1pJ&x`JT1I5Fm~`)I%L2W~9Nly)J!o^#$lhCJ9d>xLd>fvsKsHF&zx?17D8Cv`~-y77C^NAr=<99Ehi!ed|a*0Tt!!cOXv zHj5Ppi<%S9`19O{kM@i2xiKhJ+9mMXVy|{SOs*j&;Rn!ezZ?%07KQ6V7u@lol6)PH zTSk9H%_}$HJ0Y+B0Bi`mD)c4BEyqL7sE_7nMzddOehjZY4!eG$#eEqtG{ zq3C$s)dFY7NBhhly79p}VfY#JYw&9gEvWdr{5Tx8P$; zhIW4$hV^4tZt7&LI8PPMq>t{tIoIcIrTZA$N~OCA3u3pyN-z=Q6UELE%12FEId2B( zj3JVIpRI?KO`3t%lf=gm_WIafcYR3NQfuy4;q#f^V~fKkuuEPuywY+hv8EgF(f!Ei zjwsy^;O$Cx0alG&nX&C`x-82CXU#|VCkHNE)ZZ2$x`E{Kn@o8`CY>At7 zXc>7r7B`Hx*xPOTI(+7!uTH{UKY{z_b9~<31rTYndjQPoq%^Iu}yF6pt_{|Op zL*l+-;+?wjaN?##`X4iXFw@&*8)2BH=(WsQ_+p0Fwi33)JT^%~p zye8kvhGG5K$=XEvoO$;ze=!I7Xz!nQ{nTrjC)k!USljh&!SYd(^@63%?8DrNj`q9R zXeZYf$u>L(5C5=Ny9&05(hMmbJ>z#t-R1L)WSl!b`uB|fG-Yf*@4%OjdD$ea9lKhR z}~Cj z;Z;oXHd}?2VV6((GJ2ltak&HjZ}$3r7KUj`p3*-pw)tNDGT16Tdx5k^KNNK8ee~Zo z`ty{reSI5Vt70<>>%}hSr_Zg)Ss(pF3$FiBrN18@WRkJ#bYMByrTghKD&~-{+qTQ- zrz!o5@Wp4m`t`6`&YlUs`b zvc`^^{X_qjSkwHremYDaYa>jGW-DLL!Z%f4tAs6aMhxPk+$ZF-kb>~>6V6+%M#?gI zc6~SD+s}L1AZ!r39@5@;@%%se+Q)tGyY0=Rk{GLXMez3Z!K&Nok^08MR7r)(?fo}^7Vny zpHO}Cp4g^&^=Dua?5arf>}y}2U38CIzhmYgT;2y@$2tv`C3TYa_{z5!^T^lwuNr;* zj+qYl)N@|_C~W$NS<8|3>W7_vDWlHqhem&tGPbV|;PM=W&E{b>*p-u(*w^F4(SPs8 zqFQM^3r|y8nXq*1wpn}Ru;0&3FsVz%j|{Va7M^?$LBW{KYBX84`HGhLVP)9uen7nOtm(=>n1n7X(|JsE_VlqdTj{z<1!QN4!3d!@98xp)0W#T`uRGSs&d)|Iv*_RO#-A z2RY~LHgsS)*sW0C-dHy~IpphjbQ%3LrGF7V_^emI9yZIk)`NZ;{+j;R%C6Jnqy4~W zPpGl*9z3M9XJ8TRa?qAoyV}Q{C3qEo+sAqI12-nsO7~fKwzAEHrDHeGxfCF;kU&e! z`EKo{^K3)p>zKT5^tYb(*83{F=exZu4x7L(ivF(O#m<0_?ng#w=WQ6~DNE?s z+Hb|9!}Yl`3hTwLlwG5HFG)3`2vhIy8Ob(Tuy`>Qo9fFjc0N!gsgqa2sC+ba9Lr1UR3 zPW>T1|K&h(Ni*M|++nWUzD7UW$8jLrP0&^jGConqIQ7xEcTT^_E+TX@-I`<1$`rdD2X zlh<%LQrFQCX_Bv!HeG0_3tOJUyU4`ydGWr2`j^qN>%I!x&T$#d;nbV_I9$j3zVUzR z`@9-=;rct_Q?M9zdH7G4a4s};G{m^DU+~dA@?YGz)GFOW@NN2x?Qc3P6T3w?U0L31 z!%;Rkpl=@|_XcP2ae`1u`B7|b%~xT1e2>HS96yYr0Xg67@^|4?%FiiS47)tI^sz`& zELIn5h_}SAM4EW50WS+ZAMI$n+|=0CcB!t3Uf0py(bCz3sa&f0_sJ{K$ZqSF6b9?^ zUWryG_O*Y2Wj09i6=T^gC)_IYslSYnF1huP?(g&AcNq zs5bdEaCr}&ZF?I=c3?5b%WuJB_+;A-!}_r+hmWzp@;Pt*W24olv_667D6KVEDR$e8 zp)#iXj!B;8FQEA^HA?draNN18d&kkUFd4sXpDSUy-EYCDGO%L$NwSS^z`OCsES(#_I6sBJcQQup4y$Hx;#~uLJNa@`YozUbOKY19RGZYH$Jir(#yO_O z$Sl|b4Ql(>3X@~WP{+~;A7(zXZEIlj^ou^`vXF{-q@!hbF4%kCxq&xx!YDf0-v_W! z#y-2gc~~W#+Gfvw;j;7|Oj;7Z0 zP0ol<9iN&y22>r7;CZT!6<9fT@{pVyQ)`1fmH0n22YC8o{@UL5d>NLh#@T9E4=1o4 z$G2fr9avKR_D#_w0JiNAYyi6o_(=-=1llyo{9E_YsOyspn%;&ycF;p=fe_Ww+fDtP;C4 z>ftcm_2F)d)yd%%6@-IwOwRh=V>$~XCy=M=IxV)Ut}@sv{c!~D+jhln+Z7iTCaUAM zD$k*CG31#9?Zt;IU-Ob3?fE+F$EZV}!ymxpTF*Z21@RAin;rQJ_gHGxTH_FWn`=kg zHXSCvX}f6n)wT}l-cBj$+m1ojHPY@`tR_OrN70^TYu9%Zwrs{enOmw&{t!H()_Whr z$PNs#p0G6@!h`BKmSC0GrNLL(m(L{ZG9@3oS`|v`D11#_SLTQu7E1?39TvMEc9)ynBYoJm=myU=BlbS^54hY{rbAa&FX{{7Ja>f8YP=`rmENC8ocG z&yTUqRQ_+mg4k_vk4lcK{YrB)8z&77t^2ak%~85%;4`LtvaSl4{ATqKT;dXKh@X$L z{U17Kwr=_PV~;Onuv`uU&-?sdJ(oHi?~lu21ti7uwq7=2Lr1cT-GmSAf3LYgD@b z?(%dV-2uwlF?$T_R%7EPtV-FP`rl0dSN@-bOH6FdAZ!!IC(<~v$t$AcGvqj}ukciH zFvxS8Elr*LO;dTEH9F2i7ct7pyZCIqH(^U=EMu5XJi8TRyFeNtPsjRG7+HYY4U1xzi!PTZPPE-#Ms^pK1ZA0bp@1jz6LI?gUvu%v zwXC_?y5rpHE$l8yEW4cSPxsD!FTrxvm|70oCTO;QH(?SZy`H!ZkEwH@4>rO%P{T16 zQpj#o&oWWXfA}o(U*b;5JQ{gQ_c2UTjB#-+uzgwhwE?1oQqV!#Kf}wrx8fZ);~jI>!EA?r({lYft-a zFIqFkXKAbCy#B6mH4aX|wBLcxT)%tNy#5)s`dNn!SO~jKfTybDI*sY);@h>d)uZv$`bMmp(bzL)c?WjK1 z%0(oIwr3hZLzm?UWG3!7Kcs1hDj`;JnQ6Q zqCVuK|FO{@Qd$q;-HI*2Dq&$iea@{dAN|w+j~kni(s}`&rC12I!9gyQ_V^zxajMDJ z@p;?mFHy#h&ztaR#fD(=*`Yk=uvee=zxn9z-*NqqE3GHk)+n|G%ZFw8=*yUziI(>N zqS4P(S}()p`Bm`KIky_N0GlRh|3i#4aUcErMt@3a-GvV;HU*2pWWBIQpSSrek)M3d zRL1lHlK8j(|7Fc=_LJ)a2d0npEDW=hJXP;$xV%%{KGrhWDwW9RS9BAe0X-o|H$k9;iB)YWpnsqS(cpVK?nvERW~G#GB2JSmGa^1EvKdzK^EYTrJEQ7mP^Xe5v0 z5xiZo6<9eeNa9p>zI7$m)_Q?++DG^F6Vv~d#tU%$+l?WZyhC2TA1cRS?io0haP99c zqbtw;*m1i7pH*xCChvXdC8eR8e6J9@&bW{E6QeEmw)Us+h+^xoFif7grV05y;bv#Y zNAt{ocjJ(+%AbN~C{_wvLu-+g(7e#(RKau%-ZGlY*x7#GgeTX2NGI|JVT0`VkfhC` zR~lYvYT@|d)5b<@Gf47{(F`K>1@0QKM3a9h_O(UQCOx5?mawVBDCq;22$YuGv_!nK zOrDJQyQ>Aq;kJll`Z5m(9piBfe*OwZSK-~I?)#ZcenV@dqpc-U*BQCQBSRgXEp4rC zeexSonA_*Y<%ZTP%?;<{S30;)yXQRSw=Ts0-FnRl1d~`79VTLi3 literal 120792 zcmaI94_qBrU*|s{PoMiltQfKKRE%sE`=gc3UM?57RE!ugVq_~uHnNrN<BxUiH~}Yv|2zEOalVAe{~E`s$aMcc{_Lm!82`(( z!t3g z`rd#KDX$4wAG$VTn&s^=dz&Y|!QID%@_LB2Qh9B|g6MLIJ(jCYo>JwXG+u?u>o|Pp zExX(b*v3klvrN!ToX)0>_UM_ivS?>ZIe$70tRGi(ecSjiv5j}Xufs=_{}^llU7VO? zd0KgWYP|ZD*8_O9^4f$&(B%_5EH5ap6NPRL!^-O@eD7E6x`ts}f1c*75|u2kD6cn+ z*Cy+|`}i7sT6y)u#?cK)u&j^AN5*SRd3^vMJZ1YY!RpXe5}bOe<53i3dDF-H=)Z9H zS@hm|hTx00Y*ql_X@p}PmSbJbdR_N7jsH5^c>Q08w+-2B5H^Z#fZ$SL z>RuWV?+Mnee6xE#r`qX5<1Kox_kDP#s{0140$l;I&9biN@gr^?BFgJ|;aqII{z2IG zpQSl##N#}aNSiN~y4JE@`(HQyt8C--x(aXn1Dg%QhSBxy_YaBxw2%K2|I)30Lis;L zTlvej|28a$E|++m=aMq<&ctWos;zG}-lF%`{SI0kO^UuJSq_G{;|gJuh5&1!3D< z%h!ksmi4iI-FR)O>-}Z;%rDvd*aw?HH%xpv`EQDKw|1S27W8yJ8!ZhNJ7cV$QgwY` z{D+kPJ$Rk+UxQVl3lf_w&nvH^MecrADz6Ycpu7rTyYz<}M406j<@K8J(tYMd_}qlO z-#xG?d`5|Wmi4jyz<5ov-rG*zhsTxIJggC24Y9=Xo~qa3kh_l!%Ihh3zVgb1WutS5 z9G3fk&Mv>pcx5TCX87ta+S+>9;-Ap}6ET)Yl-Et;wZMAsKHh*2DX$4wAG$VTn&nC5 z^+d6|j|t`V5N)OM+J*(uKB>G0VWa2?$XE~_qRT_5lRjs_9gJ&tP zJkg@t`r|aGkmV}nb;)?`vEI9nz3_!!uvrIe2HiN3koBscx^29smDhXlgz}n&wV|sc z7Fb?YUQd^}`xsMRPr?h8R|YH>U4Y19xkqht+IZzC{vv#R+}_7p*z)^n&Kz-p)dygezjOZXLhybxd1t*9&Q9)!9tK5%_5{*PIu|Z|$ ziD5$SQ3{A1?oZ~4A)=1RBeuCuSu*!5F+MAaOk$Nw3zr=&cO&{i#P4|7J%iT%^*T7 z>--nMc4?0@zvqsd&UZ$;%1eW%sjk!Pqpyc8q91syKN~3uo)P`)KKen}DEhj``ZMLF zB}xDHSU30$c711IZRqpSf8^u03QJx?ANw_x21}CVA1vR$e*styek<=j+CR}J%Rk1t ziNkjJ5!fczoC);JK7REuxrZ2Fy&Ti-*52NZ?pK{FKH34;7}|PMc6V#6vpsq~Io^HO zM?VAWKwp5KO9|ehk2FPNuQYY=_XD48HenIA+4)^}|5Do^cMdU)Ci0 zlRo-t*eYZB5u=Zv>yfs0R`i#A^aJk0hKKe;m0)5!lC+qW(kA4+ajXndt zv-0==Cl1MY>@4%>@rJJWq#nZ z%@VARZL-)V`*E3UlPoM$B8Bna!HWI+7=Y!l?aJ?%duMkaOJ##&KKcl3^FPx5(Mx`l z``23PboppwusO7ely_4%ZKJrUlkML0(T~7}&{w0U%i>;G^ln?a?W13SHK5N))y362 z4}G-Tuprv4n_lmtHi~hc`=grISehF zj?U;8%8G(VPM_&=p7+rQVKP>kLeH?nqj#LMKKe%3Jo-M5zBK9YI9Gl2!?0oWHRxr2 zLF$9Pfchlc?rooK7GMo*lfyQW^RwqAH^t6FAMG|Qh<0l>&0$F6oIUqyw204VzIlIc za$y;0=g~?#JR9o@B|B!8Pxs^MvGv5k^O}qogEQ;y9Om^nI`iw8Mnt-spqD zP~3UPM?VGYLEA_$%yQb99=hwJTY=S}2~cl_aZcM8qrsw=9p^yh{(W>{d1%*u%gtR| zH)9XR2&F#SLf8)d>2yll>9*}=A8kEs5$%A{w#3^)!C;Yd#YZ~;8$(;4(sppI4u+kZ zKH6zm9BrY|Hor`YIcM(qXxCs>XdR<%>UR5pgWs}04*^&X+Ld?QJT!I0ljqk_A8i=6 zMc+G#R`T0a5{-ACbDSFZ+RxbIgp05pu1nINFa&mFLQj_8@bMjjjo@30FGF{yr?rLY zkoK-pmT&p!=U_4PdFZ8mcC{v6IUg;Jv*@g{F8G^vT@D;$%w*~!br>OHM3|7jpJTo5 z`|@B8Wz69v=N8AYC=x1jPWfnSU@P>u!yauJeeNY6Z2~rlHfCu{oEtvc30NQ6h@~xd z=HN9yZSUiK*Z|i(*`HMZo&K%+_iGnj0iM}}oU3mC94!cQmB~h*5BGW;hZP!MIgS;C ztbfVJyAHNQA1vbq8J~49tJQlhdZ9en6qb93*L}7bgpHz?@q)~SnQa(@hJqc%rOqVl zhv42ieJI;-%@gKrTMvBvH(?RH^9kAC-gau#Tin^ijo0C-{rSm(1<*>{mFuuOUhF!@ zZC7%hSIXnI!+8l^0-e-BuBS0V+PG{p&avs^^OBE$0yc?v>^IzNS@Dse^Olcp9M+Gn z!stTi?)c~yVfE;;jjrfOk+TNNR_)`-Z*$9)`pV#O_}jQI9$_DUihG;a89NPe{r)k& zb0IQ`$=A3p5j#K1Scj-2)&^P6f$iaQ?{)5_e}r@P&p40xT)xbhnYsElJ~N2%ALc%C ziSO*%S^hD~;_pZyI_@kgzLWcIDiZri$H}WkU%)-}H}E~-Z&0xY)*W=5el7$|%Q~A3 zB3HqEoap_>`J3{`rfI9+&3Vj4sGbH{!9l6`CdRr0><2ORL!7tIFu#yb`G3N__7CF! z``nMy2sgfhx%4{9{d-I&nOm6RcwL13FmnuSS6E3~%i%l7OZ3IW<`1ShgT%x?=X**% zSGfsad7JzEesbQ+KEJ|THIYv=u>TEoP{F6k(HL!+23PwmntP73R!^SS94G6KSw~d= zDfK(fy#v*dW5TW`Ki`4Paa|K8ZB4Fkq{NwqSM%xBKJo4DTu;KB3n`C>Xfsc{9x1jB z3!=+4V+HA#%EBd101xf^lJU(`z9-=`t+sD9Y?c0YfiTxee9QLDYe%>O^q}oi`*0n$ z^qaQl7);mc0gO%4hA79IgM08iRi`yr6}q6ePNhC|;u;;OX52-Ryq?d)9$)F>wtRe#eTQ4$3guf0Uv9DMTL{~sO|BD?Z5e#YCm^wE;fCroH`Yyrh&{Yzt;}Z#$I2%5`&wr=u zTcv!Dz;{2g&x3r}-VE13BGpc%AFqJx`o3;_x6pcz#bx-o^6i67pc^KnzTKUPlA@wg zXT-<%W8*ubeD8`@`L4if&_#$;eT(ql^6@=ZW7@y+EroCYwVj7T*bZ%Mop@AV#`OBT zyD|m_^=Kzd8<6XF8zK3)0^d{ZbO1K{>vq0AgdtCxP}1dCiZ;h6h1fPvE&CvTw$IDGpk-`_a_u|uA-b=9T|6tdl4z~A(&EA5s zY5Iyf&TqooUb3|#upxBSzmn!Ky_Tw9dl&c9g{C^uq9eseii;xAf+Q~Jem2_P-F{v^ zxe%S{Zi=_Hx15g_G{riiXL>?Rfdvym{tK5jMV@`Okh?0mIX@R~?TVg1_ewXj=$)o>1Nsus(Ecgj~~NMYyC1>s%v?se>d%R>hbXkvwoTF<#@Fb8LZd){vwQ3Y5A&tFTiKN&#rS7 zY>oD?NHDB+;_V6AR#8x91)`_Vbx7N2^C{z9Q^ve$OQ|;THoQ)iF$qhcixG0YwaQ3z zIEz01hfcct(5UbP8Cnb9B-8JIcQsKBD~FV6%)PCJ4!U z%3r*5{dh^598u|!Wa9Elc{aP!N5lWf1vr?;(q2pfIN-4E$w#J7=Hf$O?H_dRYd z>eLwL5IpdoZQpEICc0hPOtLM9>6POxEk&FsLDp&S%f?&u-uhjDPkm^6$6?dx#=PF$ zT}f}X!w}Lbvy; zkNUW#lG0Kc?l^_4)83bix9Gk5-V2}pPxiidz-G{m6Ec3F|4fuHFMSQZskR@5O_=>l zofr3EY?d~xv>(B%R9UOAYII@ZpXlRC!m5wU`q#`4kWlZw9D@ZX-Tjk#b35;>P^kiR zSuDqh4*t#}YSHQ(U56n_TQzMpwLe$k9p7u`Z5TF;u8&CNsH`~XO#1kKYJB^Z?*n+Y z^4)|*(B*r4%kgzi?XT}CnAEpG`JRBU{iU6oO4vHbX^Ak`{nF-mS80UneuIzq+s0e; z-W*?t_bcx)*Z{gX;f`a&$@{|@AKwGt>(;qP`F@JFK>6;#3en|wd}UHA6AyjtnvHL+ z@;wD#{=m*f4Qz#WGEc~{levY`k|HPO<9pNiE|_+eY7;l$3FSKh>qFOO`IaZ|!`x?% z%R?p6_Rh|8-90>leBRt9I%~2W_3-BT$$I9Df62~$22Ahwi!f;e`Ko>|zym*HZ(jvl z`(^q+LiW3-xwx$O^ef4^Pq%)hMI}^9`s_Gbd8Wi$r)#ifZWe^eHE^DV95ll$%QNt` zL>1>+K7Y&eCW9nvjTnTj5esmcn+USqGG2OLp7}TKzSL7cB#*}7Ps20+!mdjeEDPP< zb@w}&(@pK&@wRh3IZ^5qvQB$nGTx&1diTPoxmbHy2W$r2IAP9}^78I@Q8DMrppW-m z<1Kox&mDO9du+A{t4CKwFdgWeZ;N&o)2P>dd`n++^HHsQpM&rHxn1Wx(W2X;%>=}^ zlPBEGJx>MebUt1;-lF%``7(T5dH2C4&J2&u2!Kga{$dp^D=|E-&cLgjlLzP@bdp#rvXjk$b6e3`!tCg;E4IuCCd z-xajpJiG~?{YjgRz=qHzh}8JAq`N87%5`Ge$NPzw+&YWi>-`XIgYw>n1<~bdZ^pKI z{K@lLP9ENMTw}eyCt-O%@0KOknpGNif@ot|uEPO7XNd-MIuCD~x-R{KUDxaI{y(+L z8G{X=ixZE|x6&eK#>e+S)XhPU^8FNTf%4sf$upoiL~6`Z9`9*wCcjzu=sKS^-lF&B z;YIk;ZM)93u;p>ScOj(C$#XeU;0ERhZ_#_}d>TH_cT8T^09&A)OcRzj zdFJb@)bDdbLBQ*yu4$&lcmpu0VJpQSB zUC-QkXEe-D9#}WbwuJ%LPxd#DMkL<1;8Q>6=HV;(dzFKgXhnkNsiMSe0<+FzDsDm9@pXh%6AMlfG$qt;49zJcc1HMigq_$=r-So@Hnmv zE1gB3vJSPl`_QP$Itb5JWd&e4q9f#*)|DtLEJF0v?gx4tFWmqk`3gQ61p>R}0?-lr{@*RMUp&K;5$v%U*pD`cb2gY|u`QC$9 zD&I9&6}liHZBNEMMSSk~_?|fJ<{_+nkHR-`XIq4M5_1<~aaPvFgUE>g|~DS(Hr^Gn7zPx+pNul#|Xi)z>^^8yP* z>bxtDca??D&}aL2-!$H$_tyCaJfXZNV14M?Qr=`D81E{UUlZJA-IVd5OLOXpaY8)T zSQkHO`-j@yT-B&@pAoLg&4p#4+om5cV0riJZLf>Q%TZq4@R=<11mV|M2BR#>m+eV#bWHa}&1@1fJj{3JS3mtLsqaU4GL zqxSX{u#KN)KI6w9Uk8HWNb>sw&u`&nXy>${%`@YTRKIW)w)(p3`582|tYEq+xxI{8 z)`@X6x{jZ~Sd}(N+1^}z49}-7UbYUaL{~_B9UGnJVe+16^SP$(Sn@l-?xvO=zVeI4 z+RrB+=FfV0e|e{2QeKWK??w2`kX>FaY?)&_NAUQQRbFQz_2ouEZ^^Ta)qKjjW6GLR z?camE%^$P(B^Q=~PM(L8 z>rmUP&E4&>Xrwe4bRw+N$K{Ig-d64G5`0K`Ct#E4VsJTczrgd_pW`z%U&zd~GloXS ztlr~!8#XcQ)|2M+=sTZ9xUSEkPIq4$RecV^3)<~{4Zw2Ht-NNgx!ugB%6EADWGJ6? z+WUg>&QRW6@TE4}I|iGhohR^?efeTM`AP!Md6NA7eDXWqVYFE?c1I`s*G9X3E`B^9>i_lfP%zu?IGVZj>;6MX;qi9&Tx8{O99+&v=Voa*^8q58*Y+ zdl^=Xu0ng$S8VxsA3NvfOY~mvQg~j!%?e>Vw7qpga+DnZl@>cyaGj5<#&^TC)l{9Y zz!(3W?K=P)LpMlBpUm8CNl`29-^cfX@f}jW_uz5my9TR57c{;+eOz>&2DRhkdm`@U zA*_6l!gG{w7`8=QStTClp~$I$>pWaDzH6p!rSk9weEy=HhauPqx_;x!IJh{%xKj6V zWz6NyqK(mM$T5@iWR{TS`|vSSR_gjt3*UrKslMUp7uLN`jh$?}T1Hp(@;fOEf}SY=uI@H`?PLz-gpetn#s zT4H$?O%|W?Y_D_vRJXfdZK^H@;JK$aIYP~((-sqNWPn2^6@@=-pz&R zy*^LDE0uR9EE}Ce1ju{$xmYOrMUugP1+3HgxM;jZ@AdA6&)l^0(FU8P|DGUZtkWI} zh2&nW-^cd@<2$MP`}g5P-*1;S4{JnMLwq%U_B=qmgCB%2MG|&4eabo7W7@weCj`$| z^|f4~qME-!k5!_j=!ik16kISR7pgA?K@mjVy8&eS8n~y7_2Sz6aq|$~OSZ z5uNdE<@ZlriE^K7Z;2B|yTX7_>MGaR)b|5t;d!Pksaub%lWP^98+^+3jqWA+-m(N* zo#WduBx%#?KI1)j?6>UmX%^Oou8xrNDfwF|K9}Kr9k%wFFS@?<%J*q_rt-~#Wue>S zOXggb6T`MwxADqUUZ>&PU$l>318f1GX`+qgjc>KrziYhaS?@jWZ^L`uvCEx=CD6r) z8J5>C+uIySxcex-N%4Apinc&`?Z67rN8g(aCc>)oh>n>i#y_3`3fh ztLlFWzC2^szXrC#ahWHQX?z#dal8SqFSc_v0qa92zbRsOlg~|s zic6bWbmo1$pSs}IIr$qTsXQUMKLFng+1?H;58q5e&RKaTE5sZ^4!qiIm)gF|_-3K? z9?NF$knfP%yH2-v zwI`y9(n#_pUdulI&wRf*|CRsK@J!{O1~U8~@I0&mjn2{Ari>-F^B&Xd@cwB#_hYaDbaCQwTWsl+wm9qK{p3Zr{-XDKr^D;M zV6$DeFF=<~m@!IuNxY>gc82i}dR_0c##{7W@0Z}aahuh_mY8dpB~sqO?xs*V-f)s1rKWDQJ*bKUHBHLR=Y7E`gdQQq2@hRtHQ_hGg=Po><%2|Qcpo6S7jW7Z~U6wHY2c2<^@*F+(1(>`IWCz3)eZiZhTkOJi=A@ zT$i1rVc0ObK0?NIy?kd?%;%(!@2AGMU->?O$CU3TEP^iI_{w)-;_Lk2{`#JRu_moR z`JRAhD&I=jI&%d}gw@u$k7)4me%pA9-rHWT!>2oKHU=9&7bjBJ^isa=p7HTLaLKK6 zkMjK#ZN2i{ffb_5v3$!t&t$mIYlcWd9=>aQdh2=;mS^q{89F(QW_=A@ANy-=z0%iy zz}~kv;E7jkHUt|%R|}W@?dd*uCU&mdneg#_XnY5h?|pcl^4)+{pxb4PBG=LL7rq(P+&hX^&f7L32c6v4bnsb)M(_7k7^~Aa)EvVV z_%ds}WemV%o}-@n$yh1g6pP9e_hUZ34~*}S^1TO7DBm?$6*>oBse4aTb5y<<*zxf_ zG2qrWtbC8c^OSEGw#D(EJa5J^_?AV#*nXNeU&}h3i|fXFRm~|}g)e++=VBN(jIPGx zT@rn{Jw}H*<>Q_H!*0F%l=mlS+m!bfER1fOzFfY~kbE?ix#j`>cE5WlPL6MK@Rj*D z?{PW-D@fRFwi2fI_YG60O*QB68hpCe-o76;j;_IsFM7%%$?~NSvfhVUR4%tH{62qw zy*`2stNXH5n6B56%Wke}RlN?w^S{TgR}L(IZh>)woQv-7by|}5@Lj#_vG%5DPb)vO zlR8zQm1mD6S8_dQBV_pkJWG`ohb^hP-hv@dn^kiXH{l(>Y#)PZSR7pg@gd8jy>7kb z+AHr!U*U6|CZA6P`J2O)d`3LTJz*p3q+UC0r*m`SN8CDuRUMAPH?P}u2*b88UM1W< zs)V=I@Jtn4`(87?YpUJ60UuVrL$DEa{T^TGqb7WO9~$2Q<$E6 znbS5Kg$<(X@pzXT?-?KO13&84Tl8M#SQpmtGy2sus(Ecgc)O$osM@Up9h}z@qX$x zHy5J!=HdXnMtM81Jan0a+m?%i(ii5ybuPM$Zv1D1;} zKyaCGdLqn!z_e$V@yt@5&G0D(0^VHL!xrgxXH4JQ8!DF9no9+td^7EMgu=^6c%0_|pla zt)rq+SMT-yEIdz@)d*Wt_nz;-*fedCKHj^(Z^5UUZ8i?;N7q5T&+??XZwp{hE!THG zCut8OEX#Lixh(S)SJHP4y*`GaA5WH@F5?St+0Ve&zs_d4unctbXCGbPI{7VRGxkn| zb=v!i@s{z0*ZUHDKzS!%ljvf2%lUD>3zhFOtQKA7<2tvd z>b%7|8AN#NcI=9qj|$^0$Fsm_o@d?CA$wnfux*a#RC}89qOAXU+piY3%=#hL%Q52G z+Eh|f%IJ-;GV2D6x0F+9mT$pXOrKTl7@{ExvG zII+BSjledUcbH&IaM|51ZCA}hT!62=WN#OTO`{tlJoO@5|Kz(UN{ZF*bZ7C&AlXtU zX-jf1Cd>EX;+^7BMlG=kFZ^0p!zCiQ{^(D-$GB3}I|Sbuv$X}VUGDuih%n1N&)WO- zn(^9E^8***lgg_HHid4Kc$4K7b1jy6&{ZlfW64S?nags&Fk+XT93#fL#;vffo^=&0 z%NSSgK^O6;nNiP!Yu{eu8&JMy;d4J{ z*P#(MPk%c_$gyZ^eYKeH)RX63pO5z)<1KpczTAQjH`r_r7DHD{B%e>~X^FlP3MXI2 zxa#A5d{pZUqyW%%}1NL#q;hfDxGbga*mI<=Y2$# z^E`a(XYKk2VcWFVsg_6ePu}}f!*yQXbbZs-)mY#$06g^zqg$+@0gAKw>#+HD^}<$DaCt$ZV}P5Rgs(|&AB#t+(A>h0pyPn+?N;(e)A8KF84mOI;q{Hh}IlF9MAY~BJ zzFYaHtK|FsEVR1*r;Trp@_iA$S8vz97Pd^^JV&_go0s(9-{#}{uJN5m>&?sC@CoHR z2}_`h5vl7_X_3#fVcc)6pq=E}D4xX#}z4m@Iz!vDY`x_tib?3`MMPFolC*k9L+jvj2jW-|f z!AF$$EUXP(zU|#5-YY)d&yBimqfU7rf=85hHY^j}E`7V?i80*y)V+6*bvh@PjkoB% z^}YaKJ7u#tY#QB|CGHj?r?uJi14CiFK0)?ft2P zEiun9OI&1mwa#AO1Dj&qDDm}p{lDHnrr*sn#dh-Vlt@0M4sC=ie+1WaCabU>PG)cJ zkNljwAGMV0<%i*6RmU7y0G-Th9Ai24u0i>(cZUnx2oWY^yp~OjQ=t*A8Pi;A^3iPa zcNNREtSe-hsw98AD)*)fEGO{PZRR$N)oIfz|L?(Lue`?KW}!BOd3W%&COyJiK9iH_>{J z+iUPK<=YP%M>j;M=QK)6$;-Hp|0l+OQ2BohuTcK$uu61=`~Az3-~ayn{`#MSQFK~? z@;?DzyJF|161LvR_rJuWycCx?^*+8gjPEj9Z(gp!2bAw9Y!F?K$G4Ps%})FHJ~8gr zH=%qVqAgUu+pr+ITw?$DlHUy{-(=622~T}oyG%J*s+?x{!jIc|sfR5_ng1kGc`1#Q zJ8>W1TgG=5tv4?>;bSk@Y#J6v*Fc!L@Y1Gu*XgD+JlE&refSsL+=$-m{S>@Xd1u11 z(K#va;-aEh_sii3-)k1IPUqvI@fN+;yBogouWZ%^o27rBAd>Y?c~fs`_o?4K$$OHd z-H)P|e^Xd|rLU{tVl2zI;S=WGPv+7W1N1TUi5-Y~=x+)U)sS#`26CMX%Q(wB#5$Ty z^fHgoh9<&x>+GuhhN6RY`7Fm-Kf$_Pf?L7lSR{@<|Jwe2zW_s?mZ9c%y5OS+?0t{H z<`{oW5>$ovL-E!(-VfF1<9o;WPNAh54%JA0-hwy&nBB(aU@>&H#`j!2mVByX#mD!# zUv%qRr+g2=1IjlWmWghc`39+P@^?PvMNT1Hw}DH>cTdet^uniKv+LUdn?W~DkaGF$ zZJGS`cEHE?L*qN4d~d_+mG1(q0bMoW=C?RhM1I$Nd_xm%9%_{DGw|Sd+BwgKWuV)p zy-6R%y9?x%2u>KT^YEJS-BI%r7vYmXYUiN`Hid4KIDj7qt)wNMD3)i>hJC#68E?^h z+r)?P8s)tVt3_8qNS(Pam$%$Gn?Alrf62{7rSc8IcLwb`7r=Iz1K1!W$FVN=`S?n> z&cz$XcT>$7yapdrzWuOqbVG!ji=v{k?w3O? zuD0tx1{**ZCnV>mTioB*&-nNr_+_`=J<9h}wDrn&2UdtKhmi9y`R<8UzEi|Q*SFdD z<|^M)!j*3gY=wD?c|zKryaz?b7%?B;o5puR&2!vO7n^zS+w6CHU@fyS{a>CFT@n33GlI z^LzbL`MrMJ$NQG?7QMGk--M4T?`c>ZT?1iqQ7YHaMIYZozvAYiQTZN(S1I2BEJt*N zsV|4B6m@jiS=`Om`| z(bW)->dnY@)yMb9TW&sTmG5DASo!9_0_b*_Pk7|Zqe(%y&c|iryQ}6NF2L8S?0m#w z)9A(sIY;DoaB_Vf^zprGe8-jV9r%FqU4+%6t0IzZuDB>v7ViqS(B{&Am9ez(koMI; zY!l;z{GGx2>MXmSFTz-rmaocw0Ul6gSHaepKUgGcSzi7dyYIMey!1TCRruVu+T{+z zhVkhm##mne$o6_*yoOlsJy-6*HoX1^?VK#Y8qie}%PjXjX>aq)ueo0J%Ij%(rt->yWue<6SGg?fw%%>L zGL_e9_pB5XE``vg9# z`NC*K7?;2tKAFY2->LjHZ>0UCk)JH|U4-~sMe z2GGmDUtD;Q?fAQvh8Ey+2CoXXk-7Q^+snU29HX2(@!;=ymN(EWqmh4jw~mnemm&ND zls(4ZyZ8q29pQ5|fCnMJTa@ouNAYQ3Kla!knZFyP+(Kdq?|jzf|A+nU-~x=&(lS)L z?Sju@8+UsO3R|j7s1IJCMeUVjJu|caklfG19#wcs*FWgJ-RA_ zU2~edWqexH$)dC20y$7%S0lVI7uqt%YcBMU^ zi=Rn;r?}(ed*avKe1w(nQFxB3Zy2`4+`%eg#y2Irzv6VPoi@ulor~+nTlC&sUxg3# z*>xU<4WsKLXnM)9Y3MZ1IQjT~>iVYlE8hq3*f-f_ZNehxcIf9h4V-p)+YQq_o&0(( z+Ve_x^EuARxA)h*8HPMPSCw%JKK;YrkmD$0s0cq7i~0E8bbZqn)I7-z zc;fqQ-w9YBx;BDS(rN4DUkP~m)sFVAR%hPF`>Ee>kEQ6n`91*8Ro)IP4_zi9=S@#% zcW*G_{yWjRtkcJ`+jxuK>wOwNf5EPI18kv&>p#I|$LWc>Z;}qNobd6!ZM;SA^?nZ? zQ{J<%Hgt7FD%WMI`_iYpSM{DYzEypA#I!Yk(9{$NCtd-KXiT5CzS7Pc>Sw(o)=&Z=yDjdNG{xK zU$L|1;~V--Hy1U^_Zi{JHy4(HZu`XJxzgyF==q*k<+Tb?d9gwX#W_{zbY8BSayHo3 zo0lu_Q93d&8-R_W8ziLui7!N3&zFRqF(2Ot#&<~h-h;=>?XuQjRp^2QUE$;QXC`^~ zeEd(o*cJ?`Uu%Xo|4Ti=`Tu`-)Y!{X=~j5qK5r=KsDVWqR= z<9&G6&5!84-cP}^RJ}7{+2~|m!&C28@h)JU&dEjNt>;0y;TxrPd2O&+<{u^q@#Ze@+s%oyKb&3R5_*aJXKC1 zZ0B2e{~aO6mbp8ggkbyw*Lk{Xd^gZ~>wg8l7_`{{Yz*C?@s;mv#COcc_kr;pQoi@# zapk)Pt3ns__{wuEJ3hWAe$UNASot1>FRQUh7`9c({U4ERGt4!%bT`F{Te$yYozBH| z<1Kn`F0R6Rl=m=f7+s&oI}~q;$@|Twe7w`=-Fl1O>-`DZ9ObCOU7IDUhk9ejDKmfYS`*G^LtSud9Ne)QDyPuja1CX`=;?0z1RB&eDa9RCSZN& z+K2$hFL{qv=H8>t`*=U~`{w*t-Ur|rYCPq@^3Y`(@8r0|yz?!Wbvh^A##{8>dY^_* zsP`K9eBU>|Git8kU3jhXoq=_rt0z+JFWE-8W-a@8Kl7fO z57B$;{4~7cU)U@QmW6JQ`Ge0=*W~-f_&2Ky;5ruV~Szu_2j!El|AH;nHlTCeYG@B!u94;x1}L_E&3 z&%Y_mD|4sNW^ywqp7I?-8zIXd!6#K&tFYyNu*Ygg{=gi6r9BMKAphRF=D-5zc9_FB z#&XbHOQnwTJm(^xJ=`iS64^u}f6Ml@XmviW!&sHJs^&wk!iU%iZ@I&;VRU_jTr+t8 zNl2atn)LDg)cE!*-v{v6vv%$`VG(rsk9>=hzZ-RK?XT}CnAEpG`JR9;s4+n$Y@K<5 zB|_RoQ*p2;!K`z=kM9lRyR7CBuEA#t?D~$v2GR8pZXVpR#MLW7MjzjIjPIiApWlMV zl=$7rwby~&(>{9aI?8{|&iIB;P6_~ttZ-qz@gKVSo ze^-1ce_Zw1ci;oh*zIKzR*$ZVNaiUV4D-&Z@(A~98$RC8|B>r0dhhvn1il-y^NweW*NyL%>a#Dym%rXFs}DATZrEtW7h7k<$M<97JEDB=!h>qO zv;wO^7a>fWDNp8tHnZ*Hef)hl7ozv(;(7Rr8Xp8<+Z@|9!t&<&S;IP=k88$T^j_~b z;0fhD1RFuuPe`8SH(=6^CwzP#8s7osdmo;sd^cbf=n9N)Zzz;}$;+OP@5w)Qb5W># zkHZ($n5F`@K_9$KxaU02-BI5t5Sq@noa>BszS?pwS{#kFw)F7Mpy=uL zXe;x1gJ`#-_I}(oWsR#b!yR~=Dr*rY^BmcfwZgK_|8sxh*1Jx59fId7uWVQ*y2a<* zdGp8bCJcs~+4(T*wf}3ze@D$PT!haB?7Z~ArqH$F|5f-|d3l@?XAIrO7wnuoFy#!X za_+(FR5@#~Ds&FzY_cr>77K^lfh0eV-ga|Rsk}n)fbuGU?b4=Z7*~W@-nak8Ye)6B z7vXcCRr?Q{Lf3}ZA-(@CtsNb&bToB)+Qx)Wc@Ir_1FF3H@ETR#2CM?zE@KQCqxACf zq9#tBJs;nbf9mF>Q28E*?>u4eM+Izy&TSmufcUm{@x4%M35!lW>!xFNzHb=sW!1-D zgO4ijQP?25D!eCH*7<&DyatrleRz%X+JIG{+x`FC>(rSECVWi!j=+Y{RpXl+Q}U!iH{VZB!ADQqIZgjF zx4wPK_Y<^L%6AJEMz_t_fk#gsjh&*M@npZ7iPwU;mU8)fbboOSUakBiung7DTyg7? zwypZ(OYn&XyWRA z!=}-V5i(|IEej=IhA`;kd)N4mE8jct3RT}lSUtKb;w$id^c_m_ouA?4`)Id($~d;- z=A=TEQ3}sfbuWbN&`#G0lkZmYEnT8h%{rZzH;uRGz4?9}zVNio24SP<1|ECMu}peT z_;^1w-lF$<--qY@bDM3zD$o@W$>&GhW9eR0{$hW^9ZdrcVuRmNHP@ZZ?yej{w2etwFOdqKW32o{CP&odNp z`h2|a7;n*gb976z|IB7{Fq!A5CFC6CImqNY40v?UdOo!@%)ftumT4}_^Xb3joS`p} za=5fQ^F$*df9J4H?{^oBRcTqOZq4wWueR$}4_l<~mU)e2-T1d`%D9_}`}p26zO!h( z`+O5#muItSSR7pgVeY@e{QD;{`Hk9=kN4qKx6Y#XdOro<{W_av!m`mh+M9!=b6Uou zFBRzfsf75_1wDEPPd7ZrlqX!S6}3bMyy7GK7`z8#RoWbVy|;evz{j7m*%Yh?T_Z8e z@`&f`J+$8A*9)Ki9{X5#z-G{m6Unw44wjb2TgpplyMsR7ca68`z20}=b;^4Y zR*$ZVkX*>W#V+^E>ps4v|Jlt&weo!qzIW2DbDn6?$@swa4HuQ+8-eRwTycHVw^iSK z2|l5G6R=5iv0OJ7@=T-on${U%U985g$H&HdM0wwZC%)3o&kC#tT>$TOmUTT#*UbJa zujk-<-))zhCt7q{^r`MKli$R4h08c*6|B?VuN!aCdym;=`0_!U^}#044HMGNcpfk$ z3akcQgfQoYypGM;^6@?PS8gsUly50~n~Ror9|~bR^r7p7 zX$w4Y$@kxIorkN&cSH5LSKy0ZVdr50HimBSk#8}-Am{sUAKwSYcS!l(gU6Nc8mtOk z@Ud@++??+C_@21y<{_+nkHT}5Zy2^kU%X1Vc_>cip$4w=aLxFxslNLS`0#hx`5l6d zpz9}6*N4{buJRVz+9c}+sJk2wF11c&YWcso_3l&tpP+Bdvh%zJ3!}>;o_DuPuTt9` zf$zRx*DD{kM?XB9liJ^Kkz89U;Cg>wcYV{gRKI;0zWn8OefnS%=!OZp1HNPIk!SXu z5g*@=jqix^y$eq$-xXL5x`^dl>}>h?9$PoV%z2P}ms=fN=i)8nyQ2E=H{r8i zX8Vr7hR`J*`xZ-koAU8Z|H!m|<@*WR95sI0f`!rL5vd%Pc6UW$uQH~=N7wl!<1Kpc zb?_v7?yGE84O?X#u|Rmd!_Az#F(2=n##{7W?;G$&fa`sK5yq;td{w^};HxKGztq^Q3bsaHyGY1+ zZ03ud2Ke@0+S<2`?-E+?eqV?8E8j8L0J=DFi{+`$yW4yXWhCe@a|oHImv`|lNo8n+ ztNcC5X945PYCaQu24Ks4)}fL2NbIoP7V-=dP(_9NS(fkS`dBvut0o%o9AsVc-4Gl5 z^K${lCg~Y!Kf2(v-(l}_3^qp}E#s6-Hj;Z;e*aYZs{7L?XTZn*L*qZ8`u5xKdgZ?W zYe1Ky{LdQy4IlsK|JJR4mGVCV-=#5lk4-*okNJvO#xIYbeJx4;CZdvc$SwUF#(PuE zal8f}Q{MfsadZuMbBK7(z04W$@%-3$jwsK&qE((Nuo`p$JX7PU(=8lxAKzp5+;(1} zd`scm-)`?$A#8^>xK7AC<>T)~A|>)Bv>MjyTwF8$Yo_g{&Y3sh^8x!D7=n$U>nD=e zkZ`d4biAcSzPFh4@&45HP8YrRcs+oJRoypX5p+9@324?%Q**aGH}5!qx4+J(U`Wyn zlbq~j8m7pc9J7yM0c@9vo0&|j4{(3~z&f3itHyhSZM@!B;G@cW05*oM9&buvDye%p92sX2@5@P6ez1{*+EiT6`f$Q?I4`qgym@p)&#r<|w%o4XG&RnC*} z>U4X*GGMvrRuB9W<ohw6z_u8FX>V@r)yqxf=BGzH7Y4 zmG>R^K$^V|i?Dih*_L-%u^b1P(@g$UySF>q96jIK-QIdW+MMW&Mxr5k$zH~N`b&1o zyI-ZS3Uhsxx{lLvuk+a;gA_vMJI9GlI8(jJ?@9x#8ev@>+v+^M1!HyEikeG!6F%$M z`!WI>LYE-q`s23OrdA#Sb*6p1pZM>tx9Gk5@DOc-^4^98(d814&xNL#{N5lFA6@Te z<1Kox_o*beSq*H3alkwweRk@7OQ9f-yYP31kNG=t=4H+)R_x4~v< zPZNaan8xIo4)}O~XuL)5^}Y?SRNf1)26WYg^xN`GXSkf-|N8ib{)d~B8s+;8eDevr ztXx6ZhecKAs^q5jPIE8{Qw?O zzDuw=bd_FTW)e4ie4l^d=AlaY9)Yi`@jyOok2#7>LdGw>p(4IfbDvd6J%1|Cqsg(X zMZ3c_Kytw4A^C5B%kpJ-hAOKMrt^3QhCFRb&57KC5AWIM!yGJzu9lGdoHOUno#|+G zoColJu7gspoG1CjlvIX#C7*$#p>7a;#5}l>2>j3eb$JQKs+my__B-?r;g4O^xE zULd4Q_uoq^3O2DE^YOlEyhShdq$1Iqs_eCwg@-w2zh|DGb`+%*19 z!pHZv@tsENt@nHI5#>7zYeQE@WIgsx*8B9iju^;U^C=_r4{mO1R2k2}gQ|>NSO&Un z)2DiJ<3w1e^K!*_i{4xROYoKdZRaHcn?yH4JU*t8}l5V6*C){vix`+Jw^H7T;YvmkY24bk)RtmIusrQLa@3oLKpUls`{M zxjSg|u{-g<-15Szyrb~B_w4e*ur1~;RtY(F$@k+2Bh4&2wXDAP;dMep_g1Z{@$-hzeE;Y;;U*Ye-L!aXh zLy?xJ#u+E!E8GvszKDfM{qH!Kj?Iq@@mm&NNH8hw1Ug@jm!}+&YWiTfZmZwaR;s?eo!PrM!7DYNR{f)Pi>o`f9Up zseSA+-lF$Y@W70 zMWpgj5-D*KKEAh&@3fk~cn_|}EweD0=cps3zC5d1Brhjk_VIn@|GIgoSH4ffD^z{6 zU|Hz)m?Ox>_tA5A$=}k4AH5$vf==hPdt3QlAmls>@si|R!mN+~lg>XT>m67A>F`Wd?_IVpK$lI(@srxBUZ%*EU*DLQSSP!~JLdKt| z`ITa3(-(cb52d;J5WUy?Abjr=n+0Gwq9b&^TPUl5^*S#XjlZM(yWwNXzYR9am|%iP zU0>yV>G$#d!1zw8amf4dD&;#5YeZK=q|WaW`k7T9-y`X6K5CWkVR-nz+4;zU1<>s< zz7b#ENf!G;S%`mcaDN^y8{b_u7PPu8Igy-{f<~#qvfQ?)81VpZ^j!7ozvJjU(`NHBQKf?J@VTNl2aLcME-3 z^@-oK{hxcn-KRQL?jiWL>f5tnndo-8|IcH&VbJ#KHC_SbbrwFMyc%Ki_)HNUEYB#f z_l(yZ>%IAS2VSqdreHnj8i`q!2UP#?2mY|Fcp|M4F)e#6#2G`<7M_dYyNjeRy?73d0xZI&C9*YSW`$B6QJ9=@W+B|+FW zd0!)5WI4xN-#<@1a{?R}8fpih@^7Tov0TCDEPoI2Il{$Lo}Zs5aS`+xd~S1dwa)S! zNoyqfS(f|XO4cR#oMjt%)?V&;bJ=#9GIfrh_?%n+glZ2D(Kg((^R^8OqRS;@zVqew zt}`#6>pBxnxIYKV!b|&~HvT!v|3&!T-`f7Qux0w_If5!X=gWBit&itTI*iM%X<4zRYXL zxYF|+PiYZ<_xt#NVEiZ9#+#4#;X}%Q9@dDihLC!9HYK7Ze6ITV9?5j`q31yk!^6rq z2Nu9n<~6*>P|n*h>vT?DGv0b0$M@KmxjCs&zNPThkL|JwVLSBE>jam!4BMhVT_q$B;#(h8LC3HZXMT~;M*o&I)- z5MOyGVwlf*AKx3scUkq%*Wew>cN8{=uE+SwJ4=H2PW$*ik>%Dmp?n`od2iYE-G&9x z7B%yO)fw9#5QvPyUMSTMb)ftgt}H^@1mlSVp}=^6v@evQFou z+jxuK>wOwNqP!bm3-{%jmB+abhl-QeJnws4O4*M=pR(?{Wu=d^y|=78@Cr4qS%lT2 z%Z6v;mwMi(``npVTBBz=%frc6$T-_RksU*gRpJp3a0+negBjy zr-!jKbEHgPQile#JRWY<{R(WuT&Eb$Ipc)v+aP?vEDNtR%OAm6OdnJ86Cc3y)pcnJ zR);Q&x;_1sZX1ZH?Vc3QdT%~6V7cf55){@czfpd;d3-LS(e*r(?dG6S zjdc#f_f&r$faQpeIEbe`@9qBHybz5pxX@AWj{9Rh|6p%PgPctY07*+n=RBCOTLz-;pk&9s3;Kgxr5- zb0aZ9RG`h5zYp*?x;hu~q>Rl~dnRv^2($U&=uJ>HDqbx$d&fnWU{>MM><|U&1pBJv`9)xYv z2B#?Zv-t6T+@f-x39N?eoV;m#*U@`(^*VfV&SrzKQFL|qdU6t^2xr2_`=Rk3P~P|9 zHOhMfCgX)&&SR3Ar4lEa)1LmQiwwwq7aoj zR3R!=srjP=0_g~G;Rs`exKyDoaj8lj>JXJUL?sS!smeOUAu3ggLlxpum#WmEE~}+p z>SevGm-Vt<*2}8yZ&$Oo8=!mc zsEI{^(meyOQu9C=Z2iA*j=zKT4^%p7U#^gQCL5AngZ^xQ{^4&J|6@x30DODJufI$AJnYt1&7QTpjgeA#D)IULZ0xn~?N(d%7q0J3 z@DZh53tL#{nGU5bCuclB`#q~YO&M4FKDc0*E37x^;3sv$UIYSkB{e= zxYSTi@|flluTQ~q?)h03ECaiR|HNvVd;M*Vgb$2PBhg~L# zrNC>yc7<<-i=2?Ln~i_^l&wFL{Iy9s~RluIdBJ{7^!Ra^35 zJO7#Y(nwEyVn!{82l;f{Xs3$3+umF7H04u2Y!tg%_(M}JWm1)U0FVEnUuzZ?#ZLAM zc3wLr^V%{z$fpzEZ{j5Ou209{9qJs*hNWXS4{tH$QuKY03-Gmf{94tpr8VXzl8hi z8B=_eZx|jFr%#P`s@S{QAHi#tPb;uWvBRgQf56NMwW{1v_}W9iRwgVHyIp5}=hLlT zK0L^$n?^fT>|LL_;W_HOj=`p}>ww=gUT@`C(2 zdd&1+ROugr7byKSSPph8|AC+WH04qXRJk+o>3`?cz+Pp-Pv2NGTV!0lvB~Q=! zcdfSAyRo_hANhSh8-WdA7bkredQaB=YwgX+FUpw>sN?Vtn!b*yIu5`G)j6?C`8@10 zNl(rRe)IY3)aU-bZ7G}ee9CXoNbx(BG;WVyfM+qk3zuu5xi7PJ9H5U$%XAi@PW;yB zEPsz-ltp2AkF0^bI$qadJLIinSC+z>&TBpzW7utAv&E}?uataNLB6v$#NVU*Jwv^f z*n*y{<%jZ2pLeh?aCs`c=3o2S4&`P4u+BM|*x$ExHMe&)*Ech#XJN0kFI(*_rTqeY z=RLo64Q!FU!vNaiM`P(pxm5FF0lM#6-C68ipWlIxDcwm}9J`S9Ioef!F%~wn>q3C` zQ$J+l(V(;s!z+|_IxHKzW#$8k$ED83?zZ?#d`&^v%Og+6vBry6 z>Ms>@=&c24mp*OcQmM3$!wZyluGnHX%el|6^sYB?Z;aOnT*u^=)!kC-`%U=rfQgk)+zlhSO~k#e{1f+b;bF< zWxj3sDiO*+NBiDnwR4pAdErXC3bxF8JA}6MsTqgOUNk`Wp4FYl-i_5=_{0N08;A8` z7eQBSFE*7%c=2WfbPqpcVi8li2jJC8cbD>c*sZ;9-mAY{6!tQ)(=lkVn%PS80(|mU z{JvJhmRR2h(d5>VSH$lSgge`di|D*qfcCpqdzLb89NvLfEA2^G9J`RC9qDW@y?BXh zwE*p>j+>0_^F}Y#2)0B2Qd}H3P9fi%Zx9CBe zrOs=N^EL1#tZ4zd53TMb_HH~Lz=xFXEG&v$F1j*jC)R^vvmUGjXrK5I6OU@8eGDE} z+S#yl?B*BE*tl4F(kp=JINY?FyJ~OI4PW~ce;Z=3Y3w?ZnkB`2H~xOj53J^d(tHm- zs5Iwb4cO(NDf4q$dDS;nlpl<&T^&x~)q0 zAUsFedazvVR$2F@J;g*$&QV#<8c%Uu_K7k%k({PyAj zn(tXneT{J+KB6?IVKFrG(UdXYR9G77=Jn8znRa)O=f)!yo`J2)cF@SjZhhff^>^ZH zxxHM<_Luw1o`dBw2aDb)Y5IE@&xhHo6Wk15EBQNv#k=*pSg&dHG{^G$8e_PA?L|dk)V=6=Cu~zyq}~THuFD5T#n%yuQT|P zm(MXPc2*v*=;V=COC1fgY2_!d!H-q^nTJmwr(XUZCVxAOPn6qa9M56w0?rRew_2U6n>Ip#&r*Eou|I?A4P-m+q8R~SMU3HrQOrS zX@oxP<+#6?I*DB?epJ(M^4?G@_Jd#JdW|+!ohB~$HG^Nh=q}^K5_N=0n`p*q`z(IO zsH>88rxVW%d|gE&LX7G-o+frV#C(E!CKw;#r_jNjG3g81ApP1w{Acm06(7U!4*Zxd z+<#6UEJ}?1l(-$2Dd^B`-eYX@|CKpTvTaxacG)DF>|Jc_C_Kq?2^nZq+Hy(TCZn60 zqjb;1r+(M3TLqK-!w~g}ZtII}^)Y@dIU1mQ&+5)&FMcNLy9-~R_p@3t&Ty9(M`!#x)12pd$%~aW6_}UL&SDNFnUhE>oB^^!PFS~rQ_+(uCoeR)D5;k#YRoVyPWxwOE z&x7S+w@RHE=_`81XELaBG3viiI zy4qs$u`V91WIiNM$Kw^Nt?z@}fM+Z14%pOt%nM2FxLka-_0%(G3zRxRyRlKo`)~}mfyAtw*05eNu;l$W1dI2S}NyzYk>BBqn#@DZk=z#*Y5k- z1grQ4o{jt>^QQ9BEXJ-A{OR#F}GSJS#@8n!!)-`VEHYXmV z^>za4dH%;ud_t<8=iocH{q;m(a-UrjC1J*zvf_6v96>-YS+<1pDjM9}>*x+R6-#24kv z2IwCC2@{K$(menlRJyyAm;J*U{2;n^U)9|~FJ@z}W6^H4GnIA|JfyU1VX}Ydhs*eA z?7Vudv%V*>r-}z?zh||l)qdeVeEE<3@tB6m{vjXj1!zD0tcgjj(mo2$ z9r0^t!esxjNS)%h9oLDyW&wFRCO572u3FQ(;q$-l*N(xavFkwF8P{jTPi2&`#R=ZDWe+JJ|G1-7cu-m3i@msD<8)N1k zPufrI_xl12v(!AL{XD!v`CS2%{lf@6u{XKOuUGPkqq+dyx2^6X_HI1hgzM)jMq&Nf zm7|-O_d?;wOE@(hpnLE+<8NH)KE`(V5Bzm)!(^|r0Z;0-p__q*?!P9Zo06k+&%>A0 z`BDX2X5F15N&hAGm0SSJRoaK(k=Om&X|NpZR%nmJqcl`n^dhf0a2=0Mt1J5p*XI`a@@sxp2b*UN??c!6 ze6`44d-Mcozi+jtDC26s3-3|dGq6_d^7gbZ^HYL}dl<_B+Q%zQJZhBoQ}7(6odwIl zZsGUMc<#E`9WJiF?iIjwJZ@UuUDjsT=WclaKl8H~Y#O@`bZtE1C50ywd!~T^?MGI7 zOld!a=PT_6SS@y0j?bm_m%X(B-O`^iv8YtK$KkDS`+d$8TkK}xvaWFd!VAN69A34W z+iD;1GJHXeaJk##^E-6M~&kl*bsJA=*n8s)yk71p;O*OfbL_f+oyCtgO4fQ z4Oj%bZMf)0uV1N;q3c!d_xC&ulazd=dltTW*WZS6*c$8gB1y(Xa?a<_s|(P6+iHuw z8;dvLeM);2){k8WDY35dJn>0>VP`r(_u$W&zSGZ$KE`%f>2AXc(A(}nXDgpCz#~e#8n(oGJqQ<{yRW>IxCh=EpnKoy&SCGy;%)faKlQT-SPyn# z>$8lDq|fsK+DFftctn-Jo0<~c+9|BvCFf%vPPA%MlA*CKK++WEc7#?N8xJ| ze%(x1Cc2B0CKj^KV9y_tu~9&tj>S!@y~{f6_F*@?Uunl+)7W*OE%TApeyN?~fdK7C zR(nioKZMWy6MuaRuv+Z0_WYLj2s*^?^#JYDRVF4CO8W#n<5jm1z0V1S!m0>*RG58;UYGqYXQ2Y z=S(~*mF{u)(r^3Qkt?>?&Hm0_JXjCgCDvu+>3Fp(IY!WK z5%G}#?N6=tkkbALo~^W3q`cTU+Kj)%TDu*fefAejOu|b0S@`UjKPF+=ChP3P_@3Wr zC;YA^uOI*Ad`=~m+vB&a{xap)camU0t&>oGs| zThc{nygu@R{QA)7r;by`jo16|c{TRuVRhJLzxG{Mb9#PseuvYA34oJsW8 z6Z2b$yh@ns-z!#ohxORyH{g9ry8|ZAjm6NmW1D+`CEftM-|D9FT5gX&vf86c`vZ9N zm;LoE!fLR~MEmGpGxI=Hl{*CAyzMWS2Ft-NT@v-$y-H$!$|WzzuNy``HBITa!xww~ z`cc>{>v9k0**kJA4F`*P*i7A%C_Cg;<$rd&$7Dpv$w81eg> z58D|hUL@;RoS(l=#H)h5Aiv%)`l)M_NyGmXlDxhKpSa;?1F#Y7YVhj=Q!YjJ7jC)t z;34JL9IOGmDpGQcu!p)7r=C^v`fZ)=IaOL`;wAfwH1d;v9fxNszjDPEyIIcTc2h1T zQ=qXd$azb6XaKu(NE1$`sd;6J$}EcV9TuaL!8IAO}UhH z{&w4Q3m$vj&-!7b*bR`9u_`4-qwpY~J~!H_V()5y0-Smy)Z>or6yd z`~9kft&H(rBTXl2J6EvM%tSX&E97=K>Kim z>0_~XwGY5^mG&;>^RUZwv`-d?T0H(`W3T;gx7uRwYB#|{O1l;&`-f?g<99+k9-#f6 z)fRhK`#yYGoiEcc*+0~h4&W=-x}hQ##YK2Oe@ppPQn@{Ty3zK((mo2`ROd@3O!g0p z)cJxbmy)Z>or9MvpDJN1tnc$A+uru%eZeSsI+pjWw%EJvy$fGg=i)f57rWLyZMi3X ziRO880oq5J%oq@RSNkA*SB-ZMmWy46qg_&LK3kMSp7y)bYKy(A-2yLAe%Ha~S^uY; zJqFr|Iix2*`+ch|_OAB3@FjJ=&A?i*t05)l@?v|hbU8r#_(c;Fv3Iqfg2$D17Ayn1 z9oAxdE;ILqyaKq6$4#rdi><5M4WD8Wc3BKIjolc@juE+coH!>20<<4lZLxQC9>QyW z&CeEKwb)hcX*Uw94`cvD<~8Xf|!GR^^Vt zv%l&0D;t)M-8O5vZL>UIk?<=_o^JE2R$J`dHopwdQa;6Dlh}=r?AVq2A&E6?C_wvT zt1b4f_J{D9f8b|Juxjke9BuZg3GGetg5r6mB@vg@h|)g|uT}aX*cNN{B>ZJlF2z&j zI^he~{C>5oCwz-qC}!mlHjOh4=A z;ts+?$}bO=OSz0azpgU3Ca$A%$P4nT+vumJE59zo^Oau>usQY;z3{uHT*?xEyY0CH zA7NH>*$8X^yEw_lOZ-ZVziIOH_&eBY+AQ|2&SPxTJN;}MR)Aghp0@NwLOT;3t=(d^ z#opDv0H6G(pH;(_*bB^(l4D!6Wzvt4*B@{`KQQ{KGincU5598CuRjIrz^>r$^Vv32 zE~QVE`wU)554yf?z#`b?k&@$2?)Bl0m(~{S*9&m6Q}w>-9K2llRS8>RPcRIB-IPn& zPm>3z3p-7Hw(s z4tYU+aqhBdbAi%71Fus0WiZ*RjKNwVKx@c65K?JQUZx(o0oQ!XV( z`E(w>-s&$`1(Rn$=Sa!x_+osE!gc@Nv%2$YZ*Ui0OFP_t8He>^*Gfu`X<0jEOwW-Q z)ZU}*roB<6e+b_HtA717m^=fz0zYTUr4*=gXW-kMp{`$LuyxkkC6Y7l7@b_3k*C{y z*J_Kst8)jwHQ;9>umSAid)ji%W7a?Ng4%rOnrU-~(ocmCDg7NZ^08ZopE2c9BC6bJ z;i}Cc*cNN~Dk+_I%DrOxLb}4MBCnEHU)!v<*t`67`0n5Fvq9J}c72XE_dXMCmb#L^ zuPk-VV7r`R%FEtZz6)IRKZW<(d0PH%C1sGZ$)DkMi+t_lnT|x|MXl!slP{vvSxPYyPO^ z#G{C^-fgR~PM)iA3mzj-F6)PlVpoes7G?K+_lSJyGXEyATmObBD}5e0YK}jG57~M3 zr{MG7%XQ0l5li}f;5+Cm8ol=2%-3JX{(IPm9bw*}qVjac|5uoMIi4rz{ol!4#cLUp zO!@&517`uYL-{_Ke3$k(Y!&-3dRb`6_Y=r(42R(pnY_orpy;5^c~bTve5A}YHd%*x zpP#?;vB{<0Y4r2IL^;aN;8!hag*tjjy|5MRH?e8JW`VzlY0C=bb72J>Z=+e6O&;yd z#x@^KT=06uM|@c6+~2P!y39DQw&w>jiGDZ+5BHkmB+G`SW4AqM#(6itXK@ny5WLcs zOWM9-b$77E)hGAUZ@~M`_*n;R3cFDfP4Rdmv$*RT{S=`4$m)(M-G}h9e!uPlO!g`j zq{P~D?GoQgO}_`|mc~sSDwXbWc)>6Eb#uiQyG_OyA(r{>DnBK0g+s55JROVIthU&@ zv3Lc(bi>bjVdL1z*rnS=`_idXrHS{9M*_4zwc5$Co{axT@VL@mk@Dz+@N7hSDyFt^F$ZOG#FnJ}i7ur0W8--?rLf?`ppZulZ#^8->YUrGq4Wcr`9RuE^XG zpnLE|<8xf;KE}4T$FI8$E5I(BLx^7JFCcA-rE{FTiB~P(j)ocZCVR z*8{Xq-!L%|dsq7eJYQ+&!FK;P_kT$;Hrk3p(uH0bT*u=ztGkJYi_P`X>(}HfT)b-Xbxhtg`l&0FakXy48`Ru11RKJxiv7p|^b+Sw z345ym-N#n9Pw9RJPglAdun2bBoJZpKli$S%g(6I;UV86-zh8inm6EIU&%x8W{B5a( zt^5l6KGNR!3tg9(L<6+%S#7aTlS)YPdKW&?>1X4xUhG;)wm(^8LW%yI3(!9D%f@f9 zceM|~)6_iT!E&+7ARVOs#BYkXg?K7}=fB9)ec5TWQ^nrZZh;T0xugy@&px4#y-Bh! z`I;>+9-#Z4)t$!P)x8hbbJR2}hF!j+TZryrfbOxE6R}9GQ@Tgs>*_qufMsE~#2TOU zS3Z%R2iI}9VRh4#ZaaKX=|*9*e~UTV(dFZPUT=WzL#sQ9y&H!I@UYUIg+;N;wYu#+ zrKP2pdHxOFZ^u(AujTgm#NRaIqFQMmgX{N!vtjAjZTIoLEv8&*wkmf4KBnfLYSPG>Uy+vs*FJYz-3+DM0-s>l=hjsRn`eEVvTHQD#l>vk;sLtv zS>0)DU5)$jYNb03i(yym(=GBA19Xr5Efa@2rF#Uvqt4e1SQd63DV-NtpIJAWIP~(# z)3JEjYKy(=b0>UMX}7{=xRw|veGR?6dFGLaVOL4AeltIs z-(+74&@R1cVj}jg_HnqL6LQ5CyG{1y&b;0v^LiP1Iwr4KZLxQ?Ux5#(^S2i^j@=Mx z&+kwo-Xj6ppIU9PceOu)7bxu&DUV%vPh0vjam}+$9+$KU4VfR3ZFv4yOk6@r|2cS$ zwNJ6H5!eQ6^fE~_iV4UqQ$BSI=K6ON-qYb{J+KMvqRbyM*UH>=in(hDp7bYK*T+_O zMCpDA&sDlhuxjiwtZriezyiMk@3+2-zvcG$xqj2;GNoM#&%NUJxd0~nhc(g-Q!X`4 zm1~C&azDdue-t)LnI85G$#sRFv+;W2x~&hb?xb3$AHc&(cNP}KE*ISl+SAqA*m^unE31p_(SqAt`nT=)S9KhRudBUBuGnI?$y(cP%B5zia!qjk{%tL6finH*$vV)* z{g@J7JK(y_4~%Zg4EAo^?!jx6KU1&{?CMDAqMNv9V)nL+MvIbTcOPGH%%XNU#qaIm1|YFZtpurH)R2Px9+## zQOdh)4AzHTj3jZCYmZaB&IITl`a7nb9ZEM9o}qMiD4&mA7D?7T`3+;<8}%}=)&90v z-E5_M0Y1ig@A_K}TVj2kB}xDACSPIUm7+@=da(fQcdfSAyE^Z{E0p#mERLOAGss*l z+C^6i3q^c^ynbFq!#%%`4J2Zcs`nE|;OpxA&VXg1y~KHT!IVqMQ+_=UpXA(ieXW44 zvPLhE{A*hy_bP(@x@YuL=hfPN7ha|O8i(~_7s0R3Ou3XnRqkVWmhx*AR)L*7&6ezA zexK6YhU-2)`>ST`gq7~I@Fi9ax4mK5CTs8t$=!>UM!EhYPiwzvwZ-1mxeae%rgGU3 zYzVtvlFVuHyEpZh`1DV`d4GH=K>xt2#`hkj|Ao}6^tWL0W!X8T@3H!w*J6$J&CyHE z^~Is&Z|h{E*&i^TI!#@v8LF-p_;8!QEp@PY*5y9tinmO;lx4L}zX7lNF@L#Xm|R2j zkYv5+YBl$6C*eA_sejkBy;tdej%~J@qc&l34UwDFy=HXN&Y%}{ z{uV5RT@GoN<8rkwoU>X5O6v@Ky2YklNAO9jFXu`&X@=zC%NVJK!zx}kuL7o%HDn$J)mef8VsXT4^1FuN?RLm<>zEZX17O?cis}62EF5g3EWd zyZ!l!)!k9P-+=ci-456kcB7*Vcn;*XN#yYhE4SX(*HW--g6OfQT0D0b6_SkpPqV*Ern zUgI@R>ZQyQX&c)?#^Erp@?LQnLCK-6L08lee@89 zMam>=Ju)(k6R!{8LofK*EG&v$H3?~tn-+zIMa7)))E%IG;vd-lSK7zmVWpi7OUG`T z{-LSndcCnts;H0hcZB>(UR@t=S^X{5r#In)Eq>b`*aUXNr0+%N$?w}G*5SlYNBaVP)9mlYDEw{En{q`6_Pr@csX67}#q>Gqk6ctAbx*?Oia>&Emxe5B6LreQJcYDsc##+!j__g~^B4s!5z?GX5ilMTJGzIRjP#v`-A%@vpRx z!8b4Z>&u3vW4Fy-{7EcW6GLzvi&w1f4z{k(H{iocw*xkX-RKkD(vuS7fdJh{R(DM4 zK7{vQ@Yl5ftHrK@Bx6GEHPU0XT(QM&6Cb|=U3-rq&hNe@ zZZedUujBE$)!(3u8;@7v{Y`$>2OGm~kR*P~b!a$pl4oQ0>wao=hm`I|@B*c~BIU6Q zlfH&-;(lXd{wnge1GLY+Zek+#uFucHmm2*n4BKpDjFBWJd)kHdOg&x=`8pW8mwP^^c#Q_=-m|*%*t+q!3vW=mds>8*7pv)w%*SsVR7sltgc)a7BL{^19Xr6hKWN|=^hfUbkkrt z*rk(_{Z>l*bKp7-omMwP>9)Wp)E=S^HqQw$MY3@S%d=kb0NwYjuD-^&53g3b)36wt zwN_WI84HQSVu0?k+a?Zm*t+pO0^j)?ewG2t!poH3TkKu!m*GQ7I}V$~Zp5!mZQTD4(EiwJ zi@mG;Av~YvjXr7-Q=s-`*~H-^lLm zsaQ;oCi{im2CjSe)yMr&^Al71?$PkMd9RAH<6k#=sg!N7$2Z|y zcJ4Y3FJM#Ck-_@N#w^UBnd5jEF4xq(hdeKrv`E_h662MLZ62Fe(rOCV(VRdF*k*l| zxq{MU7{hw+as^RZ3UN`w$?gOl%(|?fq(2(o(R;i2Q z1&-4=&KcW3COTmhOU-ET$E^jv^}L_e!RG5dukWv#d)bYJd}VD1UjW;DkyVhk2IxPu z`jeD%=ZXjLA*DYHi(;3HJ}xtsU*spGO%JUEXrK6P<9oHzJ_Zjf?QB>&cJt@WdEG6t zJmD3BS5lYj+bdRgN3GL0;H!VxZ`%Qz!Y+ocoG0?Z@)xH#xCQDZefVqW8Ox9 z?z8{c#Gzd27QyqBZa!>>wRV~tstJGTwa+2zo+3jN;d@CVr`xL z1#`}|@%!QR9i=CWW6U$;=~&z`+9`TJ@j86&tUnfmuwk_ONKAKLr+jyIPmwnsp!6N>_+dj_7Vbjx7ttgmCPE_;R= zxQ@eHR(Dyg$8W%Ae$L;9Vb~ycl~z~gh%N^4WPon!o5tT>rTaOy4N7+t7RGMtuO!EC zNW$Qyp`rb~V0H79?(^_;rCR}8Wj`>I)D0KQV`*MpfbQE?cTw#h-h@x#gxi;+uzu{y z(UrL}G4~MP=>Xk>ziZ+USGtd}mCrJ{y4$b<>^7?HeY4`iqNa|EC51g)3!|fB(PFi; zmG%Yr?$7$Qt6@vE+|NN<#_**U8N9i>)+e(aiu!}>%lIJZu%45B1Q#k zRDkx;2@{W~(mn*QRN85<9PCzp-u0R1|HKx3*SAiio06e)Ti{z~{I+$ldDiABlEmW0 zYtck3;sLtvS>0)D-TLmsN0ja~EQVceQa5U1ycnQ+?DrCJNU2l0N8sg3Hv^W1oo98; z{8Q-V!F3#NSlu+G+YaCO8Gjq1uvymSiR50Qy*Uww-T>W)R(BFxHx3WrLrQlR7R9dG z>KcEGz2yMi<8PTb)F|Dj;1Q*p1>u)!Ft-{Y@uxaeZ zlDdh0EAsjSbU(DZ<4X5^_<+)#ht*+MX?433@h$XL19Xf2iHSp%(tR2pQo1>?H0-um zL-*oP61!IR7_x6IO3pv6?0*8ZKepN3E8QaaGG*NU%!lo;=1;R9Ncw!Lvpq3qR*IN?52ocCh9x&Ym`jc)29wr*VB zgirmHpN+!$u`9=S(M>#~7Af(j19T7GHGLUZx{tA~Rl3`-0_-+e$CJ8Yo_)(eL-%8o z)y+}5=Y=cXD%dh>^^l`m%zGKp0Ns05cOH8;4tL=b<$g8}>%}hO=!Q#Ad5^6|pVIgY zo}>0C8?Xp=+cYqXSnm6bHE$t;z07|`3^GjpGG3M$?Bd@Uc!9k}{9gWEC$p84ypPvv zY-Tu?H>;;GPUFOr_u6ZCoghWX+9I!teMk*_o!9hq_DiJvuTqv&@Fng8kfPWwQYI6f zQSpVprGAl@*Uw<`r|d6W{xQ7cC;e;{R)JlJgk<7gU@@=T0lH_WjIQi2T-|5kOF!#p zVb~^n1=%l%F87HOKaf%d*Z$tIx~tf_y061~ltHSm(7kJQm(;z6J8*e6!mVorHh^86B;(b8-%DD4o~Hcw zV7b_3kQj!)mrbKFyyB;7UTQkT>{ zhC6UQM~%P+u&Y2X;a6ekRAJ#Mc{|6OCNIdZL%*MBcWQ^yPlfN`yc^#gH1e@qui)G^ z|LGL;XTytvO(A|c72Za z$zt>Q(1`%;$5vbHU7gS1IZAs27Qs&LYnXGk$o#0$^nRbu!)Q)QzS2DlpI7^ha@ZQ{ zw%phF3WdyjJ0)IifbLzZyM(Rl?;Utd>5jk#u#@{5Pjs2{rvh{j{DJYeN9lfnZI;sA zf`zb?`x+7ld7fP6=X5l5-@RaUbCvEn`0R83HdMk^Sf}SnMz^qtt%lbSp!<&1UBK3j z!&~rbfXl{Ueb~jUuFSQC=+40Vc@c$x8V)S zry|O28u}@PzZNkFXZ9VIGr%kz(3RUg|e9`tt za?h0q+hyI{Ano~7?3KX@g|FV%tnQ{-XJ3KGm2EF<9J?V$H(Xli4a3=n`E)i;!eY6*FN-Y*8=GPzV);jJ6seK z-{aV?sdM-}Ti-MquAleeBZ^JKVz64$9LHNq>&UE$L#rxx5T3(eNyCTa*!Ez#*kwqP zaw`5!R!i?E&I?!iRj_6HWR6tJ@z_CAmyE$tP8_NC7JR+j9Dg@|*GQ(!{bbI4+4FL- ztKoPU2e)~xBcSrEas_D_(*bOA$?M>Fp4Tj1ThSQ8E|a_%WtOQ!_vi6{k%&oZ4ZdOW zr2S99)4$SK~eS!YMzSgEe56gD%bY_(30UIY9IHecOJe`4l{^G_znC*hyV7p1UIEM{^6{ zItDkb?(R?d{pyD2DBT!r8oLg-_EaE%h6TQq0lP^EOAhR2j{4lE73IXI`X*WK8D z<)v%&GQ_+Hc{&!ito9abz8lY*@JyxM1Dn7u3YUKC!uP=d&5x|+sM7oZ-jVIEXAxF| zoz#_%X45z83qz5T#GJbkp#AJ0n^=@9?IL&s=e*m7eAo{AhG}?W?DEh|8BE9EHLJOa zty}&Tc!tvKg^gn;bx90*@Ovmg^JA+yqBK8*$G*p3&l0Q}JE=?hELIXKZb>}j&vPFE z+RwdfVo;{EOW}1Te(eI7TtiI3524+0H6AW1j?{-QnfalLd>xNFR)1BkE3d-PKLd)0L{;>=77@t1YUQtwFD zNEO~xfbM~L<8zPF{Q}!6rMm?SVYkUnFcITZCEhX0x7xCD-V8Aq!zB4UTQ4bxGz)Ja ziH~9Ib=zJwKBR1`J@?D-g)D#D;;>2VlGh9`wp~f==>`LIKeD=`O7{bJhtgey)nHd{ zZM&MU$~M7U570gRuM#m!sZhEn;Mqzy54Owtx~ADDQAmF~0f8hib~FfjcO+hm1VA&C#MMjkZcUlm-(;*Ql_wYEu(*Wp9o>2Jdz zY#6&flK9)!9&hg9bv!`#3#;3&bU%fUvp%`8S%;Njmv3#GzS${ZNZeOUe5Y=t%5DHj?Iy3Io7ebYqgfhck8$VuPyMi5!e8BvNw@2 zXTSTGzgU9;v=1&AKlNVbF}7t))^2^3)2d;iMd`RiXV6uN0BRRe&_Lc(y z+K;TZ*t^;f;d4LgXA7`e>?%lmzL((p#f1Jwfc~@Znb?#o{UW%0p3|)}AGTA%TuR#W zy(FPsL7tAwZL2NzuFfs^j2e&quu<#={MvGx(;ExW{@iMdy{r8RyqZPcWoxi<>pcB$O?+yV{!w^_($9ot zVz@Gey#|C@C!RddY84zCN`2lWJ@}5ZiRW{wyqt zT{UTeW9{oxOQzlWJ7x~U<@*5KSf|6XrB2d5U-=YcKKWYzWvj2hW2O^6{j^`d6*lu@ z?Bz)N^&{SZlu_sQBdb3~8Q0f`aCwfxWec!c>?%mh9P4@F_`fr;s8L!^!PAvi7Ayn1 z9rhl%91n04OzM*PBhwzAhbO-WA`?E%rf2}84!Nkpa6=fc>HIC9z3Am{`ELAdUbA_f z#Ws&JW909U2GG;-eFPJ=F|`MJ2%q~Q)qYqlc3H%ouJzilU1@8qFRt(Acs)S-^h2Xv zp|nrHqe?puw#$BF2F{e1)IL#If6*%^Py7A4)!tBhlvm+XwAJmmKG+y`b!dMbZTr2n z@uvEol2A#HHx{7(xz!&~`k%lvf55N51}n#IH_z-_63=FxD!E8i-fn>QxqolsQlPZY zz_;l~*Y`5mI{Se!w7-mANhlQV;W#ovpMlz#~e#8n(ooHVBtBFcz=xC=9iE ztpU3Ct?rzf1K)2G|_)S}*$P_-n^sJBMCxfc8VHJ*nox2k@}co`prR z%SBsaZM08%%kXOccH4RU0~3=PrTY{-N7-hTf^eulHqm@Avsx95#tvEBbqWmw1B#x*u8HQKkC=XiqmRx~2co_*SVjj>C0)a>W+AP1e;!e7H9dhE?0L!a7LO_RH|}Jb!x| zU^+JMh$eCiXu2_b3$D+VF<2jVWwh-}6uJ_Lc%Pf{DFdqfCvffm8mt_<0#d?%uJuKy zn9;jSp5CXO_>ldM)nyChO&??bLfYh<8X)D7=5pzG(hlQxAj|XC=rp-*l|5Nm3UmKK zV)C67PUpJ#tE{t3taDTyqiocpAqjXsec76x9%ITom|7{4zI!H$Kg8O?^*w+f6%XSAFe+WJ`Iau zmyiDp3GY&KXJd>D`^5m=WBJW4tYXJ=)CI z8t`SI7h5~8U2SfUkjJ-!E?+{n9It-KpZ|r`>Q`Ey!gH0@I;;%49p+G((*x%uU-uW% z{a2c#`wMAz-rD|o^gK-FZ`bE4nC|yGFse*lwc{k&$FIYC@W(Ab2ph((8lLF)D=&4n zUG==5Mc~_+qxPm%J2~-P1AIF9YIA$^+Qqi^Ycf%7QHP#mkNs!PF*QeK!xrgK*S~g{ zv?Wuur3pU5dgR*H!WI}8{j6nS74v9k+umBR|GYDSw`;-}I@;fdurcO7x4s2f6@%Jk zPkm(kZLs4r8T-TV-S76>ro*zaTZaGiw;AumnoztqYvk{(33&gESBIwd{T6Ip&1d~E z9qSKa66gQE8pi~b=YOYCDz|^tOP8;ogK0I>xmVNy~qEEZCwqvKqH< zh$aDWZHHll*j2($Q|K@5Pv1$DA7{sv*i^AESrcE-aP<%W7xqtTj!cK?{(J!@KIbZ* z&%tNa`CbWIVVn$G-qiV0BV|3d+uw%&do$OQ4I#ddqPe8zp!Z?pcHWdZvY%9j%_gr4 zoZvD(Yk6J4PPga8$EH1cKXD8$`*k-)*|2olGY_YG6Ys^x>Sbs9qU<)3!}Xf%b1Sgb zb=|ggt*fzd3!bLx>W7VDR|}Wn%I|4)IyKZ6hU@b$HWk;0i_K@ArUU9b_&=Gp#Z`Te zu^rCzw{06%fZYb1?ft&`ItvS>!jgJE2o*1=zj(4Z5~n-70_xE1ylLy*RpY4}o}=oG z!KSh6fOATEU5!_}uD4%un_3L^h6C#P#MU#c>UkuZTwl3;w+ySoE}eR4rr96v^;o?$ zuedN0k~Vo8``dILMowzJs_U%Ss=CTyYmCQHcwpa^n7%71E=p9#ZB?E_;bO?Q2)dgO z*}mo{2ii0B*pE|(K8HVm$+e!_?nUtrdzT&ipG{lp)L!EVe1~gC*ERztziGQ<`HgFx zGQ7Pq(zl$0oNJ`Lvsf*Jly60Qj#sz7+prZo_sQB)WAlgMQMKRu1V(o1F#8Et;}JZh z+OZ6)!Y&=Y#<9F7;g%`=#AsD2t>f@@bzPY&w%EvuRySAao`cWY^2xd?Ve*^R!*Gd9Yh(OkE3f}e=Zsuz@oyNd zU0&U|KY&fzxmV)RWAo?W+V3Nuns~ITF?tZ5ru_C`GL}~vyE$kEtV!m}$h*z_9l~DM z_lm7=M~$@`@L_dMbik&tllS%MUhmSiYjX3cqphboCbug<-|jUk!+&pkx`6H=W!;#4 z0qaq7;})!1*`4{{?D$vypN30JT+I+{3&$tZIkCwrX5cgBcME5_?0X_P!2>(5|hr4HJ2BRT#)hRbiDyV|R; z3JEftF80N`zKa`1&A_ewb)1Ju9r9c4uHIR=d>5h1%3*6XRzBauY2vFRar^1gH7^=a z$30ueydBfYcHD(Gs5-`Bz1T(Ivghq=yn6LIK@M^JZOoH=B_MY)WAzLcbAD~3w@F#I z9Z&!7^i76g(yxRp@!m!*T=zjYjO^5O<@;rLzPj#dfXy)m<$6fw1@1XDO8@o*=)P}t zr_`8z7rsPc*WVdfD|UJ4a(UvlcDUQfp5oGwZ1XM^@??G@PXGUVF5bD8J=e9KIJbI> zdP)=9E-&Y^{cGRzusk)VR={=$n(NKy2YjdBju(#9eRxqa%}CaU!h zpJn|^+$mW=BVXx$0h8aK5WkXr+Q;#>_`s{{@7e#u_#3w8qS)tKo9D<|#Mrfuz&02! z07SV_JLgqYR>llJ;n2|JT<3R!1OtC8%8y$ z>uRjrg7>NYXg_QeJNXQ_oZr!A`ETOyj)f=NHQzYMzuaZd>qT^WDK9Zh&L@vyaZYBJ zZNs#`Xa7&z|H|KI;rbdp43quC1i_WQy>>Bvt%C{a1jm1|zol}nJsG!sXw6!mrLU6f z`a8naJU9u{ey4tJ{O(oj`sdi{XB{?SVeGaD4qfQAv~@Puw?+$P9zAPyB;SqIvvB!c z7?*`%n+Qyx! zuh>+pa!39z6O&d|$3eJ0S3FoQcB^Pe{EJG0>*5#KyaH@>T{mrAyK0Pe2hB-keK{{&u`1XyV|I>2NX@hV@km6*S8IUweXA|@uJ+sTam6NJ zJ+P=>n{Ni-9#w$$(f`}{E%vVVA$X-?X|NoaJO}9bo%rI1T(PGPSEJi%>*oe9!#DN% z51WHclJ>MuCj9OV(0*vO#on#+0bD-g?6O&ye2;UrPrI~;Vdbp^XrK7P#6;{}onvtM z-e{L)!{mE2w^=y${4On$>%1^|Iwr4LZLxQCUWTtI7Kcs3Mo4U-eEZ}QqCOm;|B2Nf zR$7nXJ&G;Es$h{Iea@}z0R6N7kBLoKX*~(G6eg20^y&Cd$eBQG9%an2B z^9FoIv0<3JcPP&}?APb}-vaax?i&B&O6xJUwTf-S3Sij*`ZA|xp{4!5Z1uC0)(dcX zeii&=&8>kg!e&U?|1dL6JV5_Ft3R!@?!!kEn})?;vR~Mz&$sz3lb?LfROa+SlK6M< z|7Fi?kCXcY52o9C9)?*;zN+^uT)tD?ZEHDfjY{PGD~1WrfS!_{I!UZo@s=20k9;E9 z+}(Dux&G=k-lz9&VZV#5XfWM)`BD~VvZxt^E?`bb#*J$9DWHjc4Kdw;RJS`3`yce5ka+-ZSv3;M(6iR#%?=apU$n zd`_`Jn0)U;A1NK(pGajz|sHW>-!^n^_nMoAyIM4+_frYGW^ZS!Ql-`g#C z9=An2JC=Dk=md{r@bgzNh6>+ZYQCSz<~O!SJFm4x>${>?cx0%vtL<95sZV|*3Ul|o zxZ2o$y`}MD{CX$%Y4@GS{MLo|zgMq4fnXBr;uD35XX$PJwT@;UVybWM+N;&`{y(}a B6$AhP From fc34c824d349c4ee97b9706a29e388b90b0573ba Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Tue, 1 Sep 2026 12:32:21 -0700 Subject: [PATCH 33/43] Automated Code Change PiperOrigin-RevId: 974645808 --- third_party/xla/xla/tsl/lib/gtl/value_or_die.cc | 4 ++-- third_party/xla/xla/tsl/lib/gtl/value_or_die.h | 5 ++--- 2 files changed, 4 insertions(+), 5 deletions(-) diff --git a/third_party/xla/xla/tsl/lib/gtl/value_or_die.cc b/third_party/xla/xla/tsl/lib/gtl/value_or_die.cc index 87d053c974d660..c64c1f71f1f31d 100644 --- a/third_party/xla/xla/tsl/lib/gtl/value_or_die.cc +++ b/third_party/xla/xla/tsl/lib/gtl/value_or_die.cc @@ -21,8 +21,8 @@ limitations under the License. namespace tsl::gtl::internal_value_or_die { -ABSL_ATTRIBUTE_NORETURN void DieBecauseEmptyValue(const char* file, int line, - const absl::Status* status) { +[[noreturn]] void DieBecauseEmptyValue(const char* file, int line, + const absl::Status* status) { if (status == nullptr) { LOG(FATAL).AtLocation(file, line) << "ValueOrDie on empty value."; } else { diff --git a/third_party/xla/xla/tsl/lib/gtl/value_or_die.h b/third_party/xla/xla/tsl/lib/gtl/value_or_die.h index dd7ade6a99c348..6acebddefc2630 100644 --- a/third_party/xla/xla/tsl/lib/gtl/value_or_die.h +++ b/third_party/xla/xla/tsl/lib/gtl/value_or_die.h @@ -37,9 +37,8 @@ namespace gtl { namespace internal_value_or_die { // LOG(FATAL), with a source location and an optional 'status' for details. -ABSL_ATTRIBUTE_NORETURN -void DieBecauseEmptyValue(const char* file, int line, - const absl::Status* status = nullptr); +[[noreturn]] void DieBecauseEmptyValue(const char* file, int line, + const absl::Status* status = nullptr); // SFINAE helper to detect instances of StatusOr. template From 22d32c4a0e13b2533403f3a2bdd6956ce54d6682 Mon Sep 17 00:00:00 2001 From: Bhatu Date: Tue, 1 Sep 2026 12:39:10 -0700 Subject: [PATCH 34/43] Support tuple entry parameters in MakeDataflowConstrainedArguments. Return the full instruction constraint map from ConstraintPropagator::Run and unpack tuple entry parameters in MakeDataflowConstrainedArguments so that GetTupleElement user constraints are propagated to individual tuple elements during test argument generation. PiperOrigin-RevId: 974649257 --- .../xla/xla/tests/constraint_propagator.cc | 8 +- third_party/xla/xla/tests/test_utils.cc | 88 +++++++++++++------ third_party/xla/xla/tests/test_utils_test.cc | 33 +++++++ 3 files changed, 94 insertions(+), 35 deletions(-) diff --git a/third_party/xla/xla/tests/constraint_propagator.cc b/third_party/xla/xla/tests/constraint_propagator.cc index 70b33e7bd1000a..b504488f97fbbe 100644 --- a/third_party/xla/xla/tests/constraint_propagator.cc +++ b/third_party/xla/xla/tests/constraint_propagator.cc @@ -394,13 +394,7 @@ ConstraintPropagator::Run( } } while (before != propagator.states_); - // Extract only the parameters - absl::flat_hash_map result; - for (const HloInstruction* param : - module.entry_computation()->parameter_instructions()) { - result[param] = propagator.states_[param]; - } - return result; + return propagator.states_; } // Accurately modeling full relational semantics across multi-branch graphs diff --git a/third_party/xla/xla/tests/test_utils.cc b/third_party/xla/xla/tests/test_utils.cc index da5540e06f576c..49135ab69658a7 100644 --- a/third_party/xla/xla/tests/test_utils.cc +++ b/third_party/xla/xla/tests/test_utils.cc @@ -422,31 +422,17 @@ absl::StatusOr> MakeDataflowConstrainedArguments( ABSL_ASSIGN_OR_RETURN(auto constraint_states, ConstraintPropagator::Run(*module, get_index_known_zeroes)); - const auto params = module->entry_computation()->parameter_instructions(); - std::vector arguments(params.size()); - for (int i = 0; i < params.size(); ++i) { - const HloModuleConfig& module_config = module->config(); - const Shape& param_shape = (module_config.has_entry_computation_layout() && - module_config.entry_computation_layout() - .parameter_layout(i) - .shape() - .is_static()) - ? module_config.entry_computation_layout() - .parameter_layout(i) - .shape() - : params[i]->shape(); - - const ConstraintState& state = constraint_states[params[i]]; + auto make_literal_for_state = + [&](const Shape& shape, const ConstraintState& state, + absl::string_view target_name) -> absl::StatusOr { ConstraintInterval interval = state.GetConstraintInterval(); StructuralConstraints structure = state.GetStructuralConstraints(); - if (!generate_aligned_ds_indices) { structure.alignment = std::nullopt; } - std::optional> limit = std::nullopt; - if (ShapeUtil::ElementIsIntegral(param_shape) && - !interval.IsUnconstrained() && !interval.IsEmpty()) { + if (ShapeUtil::ElementIsIntegral(shape) && !interval.IsUnconstrained() && + !interval.IsEmpty()) { // Use exact hexadecimal floating-point literals 0x1.0p63 (2^63) and // -0x1.0p63 (-2^63) for boundary comparisons. INT64_MAX (2^63 - 1) // cannot be exactly represented in a 53-bit mantissa double and rounds @@ -476,20 +462,66 @@ absl::StatusOr> MakeDataflowConstrainedArguments( "Unsatisfiable integer constraint interval [%f, %f]%s for " "parameter %s: collapsed to empty discrete range [%d, %d].", interval.min, interval.max, - interval.exclude_zero ? " (excl 0)" : "", params[i]->name(), - min_val, max_val); + interval.exclude_zero ? " (excl 0)" : "", target_name, min_val, + max_val); } limit = {min_val, max_val}; } + return MakeFakeLiteral(shape, engine, limit, structure.needs_sorted_indices, + structure.no_duplicates, use_large_range, + max_bits_of_precision, structure.alignment, + structure.known_zeroes_mask, + /*float_generator=*/nullptr, interval); + }; - ABSL_ASSIGN_OR_RETURN( - arguments[i], - MakeFakeLiteral(param_shape, engine, limit, - structure.needs_sorted_indices, structure.no_duplicates, - use_large_range, max_bits_of_precision, - structure.alignment, structure.known_zeroes_mask, - /*float_generator=*/nullptr, interval)); + const auto params = module->entry_computation()->parameter_instructions(); + const HloModuleConfig& module_config = module->config(); + auto get_param_shape = [&](int i) -> const Shape& { + if (module_config.has_entry_computation_layout() && + module_config.entry_computation_layout() + .parameter_layout(i) + .shape() + .is_static()) { + return module_config.entry_computation_layout() + .parameter_layout(i) + .shape(); + } + return params[i]->shape(); + }; + + std::vector arguments(params.size()); + for (int i = 0; i < params.size(); ++i) { + const Shape& param_shape = get_param_shape(i); + + if (param_shape.IsTuple()) { + std::vector elements; + elements.reserve(param_shape.tuple_shapes().size()); + for (int64_t j = 0; j < param_shape.tuple_shapes().size(); ++j) { + ConstraintState elem_state; + for (const HloInstruction* user : params[i]->users()) { + if (user->opcode() == HloOpcode::kGetTupleElement && + user->tuple_index() == j) { + auto it = constraint_states.find(user); + if (it != constraint_states.end()) { + elem_state.Merge(it->second); + } + } + } + ABSL_ASSIGN_OR_RETURN( + Literal elem_lit, + make_literal_for_state( + param_shape.tuple_shapes(j), elem_state, + absl::StrFormat("%s (element %d)", params[i]->name(), j))); + elements.push_back(std::move(elem_lit)); + } + arguments[i] = LiteralUtil::MakeTupleOwned(std::move(elements)); + } else { + ABSL_ASSIGN_OR_RETURN( + arguments[i], + make_literal_for_state(param_shape, constraint_states[params[i]], + params[i]->name())); + } } return std::move(arguments); } diff --git a/third_party/xla/xla/tests/test_utils_test.cc b/third_party/xla/xla/tests/test_utils_test.cc index 638b39a83d8ee6..6c331b7ae30b00 100644 --- a/third_party/xla/xla/tests/test_utils_test.cc +++ b/third_party/xla/xla/tests/test_utils_test.cc @@ -274,6 +274,39 @@ ENTRY cluster_13361217111314620287__.11 { } } +TEST_F(TestUtilsTest, MakeDataflowConstrainedArgumentsForTupleParam) { + auto module = ParseAndReturnVerifiedModule(R"( +HloModule cluster_tuple_gather, entry_computation_layout={((s32[10], bf16[100,256]))->(bf16[10,256])} + +ENTRY cluster { + arg_tuple.1 = (s32[10], bf16[100,256]) parameter(0) + get-tuple-element.0 = s32[10] get-tuple-element(arg_tuple.1), index=0 + get-tuple-element.1 = bf16[100,256] get-tuple-element(arg_tuple.1), index=1 + gather.2 = bf16[10,256] gather(get-tuple-element.1, get-tuple-element.0), + offset_dims={1}, collapsed_slice_dims={0}, start_index_map={0}, + index_vector_dim=1, slice_sizes={1,256} + ROOT tuple.3 = (bf16[10,256]) tuple(gather.2) +} +)") + .value(); + + TF_ASSERT_OK_AND_ASSIGN(std::vector args, + MakeDataflowConstrainedArguments(module.get())); + ASSERT_EQ(args.size(), 1); + ASSERT_TRUE(args[0].shape().IsTuple()); + ASSERT_EQ(args[0].shape().tuple_shapes().size(), 2); + + const Shape& indices_shape = args[0].shape().tuple_shapes()[0]; + EXPECT_TRUE(ShapeUtil::Equal(indices_shape, ShapeUtil::MakeShape(S32, {10}))) + << ShapeUtil::HumanString(indices_shape); + const std::vector results = args[0].DecomposeTuple(); + auto indices = results[0].data(); + for (const auto index : indices) { + EXPECT_GE(index, 0); + EXPECT_LE(index, 99); + } +} + TEST_F(TestUtilsTest, MakeFakeArgumentsForScatter) { auto module = ParseAndReturnVerifiedModule(R"( HloModule Test From 5853a3b987a85ef909d53bce2d699248d2822162 Mon Sep 17 00:00:00 2001 From: Maxime France-Pillois Date: Tue, 1 Sep 2026 12:47:02 -0700 Subject: [PATCH 35/43] PR #47769: Reject transpose tiles exceeding shared memory during Triton tile selection MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit Imported from GitHub PR https://github.com/openxla/xla/pull/47769 📝 Summary of Changes Add a shared-memory constraint to the Triton tile-selection search for transpose ops so that tiles requiring more shared memory than the device provides are rejected before compilation, letting the search fall back to a smaller tile instead of failing later with a RESOURCE_EXHAUSTED error. A transpose stages its (padded) operand tile in shared memory to perform the layout conversion. The estimate is `product(power-of-2-padded operand tile sizes) * element_byte_size`, compared against `shared_memory_per_block_optin()`. 🎯 Justification The Triton tile selection process did not take into account the amount of shared memory required when choosing the tile size. As a result, some merged kernels failed with a RESOURCE_EXHAUSTED error. Example of kernel failing on gfx1201: ``` HloModule transpose_shared_memory_repro, entry_computation_layout={(bf16[1024,2080]{0,1})->bf16[1024,16,128]{2,1,0}} transpose_fusion { param_0 = bf16[1024,2080]{0,1} parameter(0) bitcast_0 = bf16[2080,1024]{1,0} bitcast(param_0) slice = bf16[2048,1024]{1,0} slice(bitcast_0), slice={[32:2080], [0:1024]} transpose = bf16[1024,2048]{1,0} transpose(slice), dimensions={1,0} ROOT bitcast_1 = bf16[1024,16,128]{2,1,0} bitcast(transpose) } ENTRY main { param_0 = bf16[1024,2080]{0,1} parameter(0) ROOT fusion = bf16[1024,16,128]{2,1,0} fusion(param_0), kind=kCustom, calls=transpose_fusion, backend_config={"fusion_backend_config":{"kind":"__triton","block_level_fusion_config":{"output_tiles":[{"sizes":["32","16","128"]}],"num_warps":"8","num_ctas":1,"num_stages":1,"is_tma_allowed":false,"is_warp_specialization_allowed":false,"waves_per_eu":0}}} } ``` ``` RESOURCE_EXHAUSTED: Shared memory size limit exceeded: requested 131072, available: 65536, context: [Fusion: fusion = bf16[1024,16,128]{2,1,0} fusion(param_0.1), kind=kCustom, calls=transpose_fusion, backend_config={"fusion_backend_config":{"kind":"__triton","block_level_fusion_config":{"output_tiles":[{"sizes":["32","16","128"]}],"num_warps":"8","num_ctas":1,"num_stages":1,"is_tma_allowed":false,"is_warp_specialization_allowed":false,"waves_per_eu":0}}}Computation: transpose_fusion { param_0 = bf16[1024,2080]{0,1} parameter(0) bitcast_0 = bf16[2080,1024]{1,0} bitcast(param_0) slice = bf16[2048,1024]{1,0} slice(bitcast_0), slice={[32:2080], [0:1024]} transpose = bf16[1024,2048]{1,0} transpose(slice), dimensions={1,0} ROOT bitcast_1 = bf16[1024,16,128]{2,1,0} bitcast(transpose) ``` the tile selection must therefore take into consideration the use of shared memory when selecting the tile sizes to reject incompatible tile sizes cleanly instead of failing with a RESOURCE_EXHAUSTED error. 🚀 Kind of Contribution Please remove what does not apply: 🐛 Bug Fix, 🧪 Unit Tests: New tests included in this PR Copybara import of the project: -- 284e4ae61b7b9b79b209747d597ce5c39163bf1b by Maxime France-Pillois : Reject transpose tiles exceeding shared memory during Triton tile selection Add a shared-memory constraint to the Triton tile-selection search for transpose ops so that tiles requiring more shared memory than the device provides are rejected before compilation, letting the search fall back to a smaller tile instead of failing later with a RESOURCE_EXHAUSTED error. A transpose stages its (padded) operand tile in shared memory to perform the layout conversion. The estimate is product(power-of-2-padded operand tile sizes) * element_byte_size, compared against shared_memory_per_block_optin(). Because the tile search uses different constraint filters per path, the check is added to both: - TritonEmitterConstraints::ParametersSatisfyConstraints (legacy/symbolic tiling path). - experimental::VerifyTritonConstraints (experimental tiling path). Add a shared GetPaddedTileSizeInBytes helper and tests covering both paths. -- e066b323035ce1da9210e252b29e3345bbaf8e0a by Maxime France-Pillois : Remove unnecessacy tile size check Merging this change closes #47769 PiperOrigin-RevId: 974653326 --- .../gpu/model/triton_emitter_constraints.cc | 86 +++++++++++- .../gpu/model/triton_emitter_constraints.h | 13 ++ .../model/triton_emitter_constraints_test.cc | 128 ++++++++++++++++++ 3 files changed, 226 insertions(+), 1 deletion(-) diff --git a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.cc b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.cc index 19746e7f47456b..ff4683114da0fc 100644 --- a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.cc +++ b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.cc @@ -48,6 +48,7 @@ limitations under the License. #include "xla/hlo/utils/hlo_traversal.h" #include "xla/service/decision.h" #include "xla/shape.h" +#include "xla/shape_util.h" #include "xla/stream_executor/device_description.h" #include "xla/util.h" @@ -75,6 +76,18 @@ llvm::SmallVector GetPaddedTileSizes( return result; } +// Returns a conservative estimate (in bytes) of the memory required to stage a +// tile whose sizes are `tile_sizes` for an instruction whose element type +// occupies `element_byte_size` bytes. +int64_t GetPaddedTileSizeInBytes(absl::Span tile_sizes, + int64_t element_byte_size) { + int64_t padded_tile_elements = 1; + for (int64_t size : tile_sizes) { + padded_tile_elements *= llvm::PowerOf2Ceil(size); + } + return padded_tile_elements * element_byte_size; +} + } // namespace /*static*/ std::vector @@ -120,6 +133,7 @@ TritonEmitterConstraints::GetBuilder( const HloFusionAdaptor& fusion_adaptor) { absl::flat_hash_set unique_tile_size_maps; llvm::SmallVector root_infos; + llvm::SmallVector transpose_infos; auto roots = fusion_adaptor.GetRoots(); for (const auto& tiled_hlo_instruction : instructions) { unique_tile_size_maps.insert( @@ -134,6 +148,17 @@ TritonEmitterConstraints::GetBuilder( shape.IsArray() ? SpanToVector(shape.dimensions()) : std::vector()}); } + // A transpose stages its operand tile in shared memory to perform the + // layout conversion, so record the info needed to estimate that usage. + const HloInstruction* hlo = tiled_hlo_instruction->hlo(); + if (hlo->opcode() == HloOpcode::kTranspose && + fusion_adaptor.ContainsInstruction(hlo)) { + const auto& operand = tiled_hlo_instruction->operands().front(); + transpose_infos.push_back( + TransposeTileInfo{operand->symbolic_tile().size_map(), + ShapeUtil::ByteSizeOfPrimitiveType( + operand->hlo()->shape().element_type())}); + } } std::vector custom_constraints = @@ -148,7 +173,7 @@ TritonEmitterConstraints::GetBuilder( return std::unique_ptr( absl::WrapUnique(new TritonEmitterConstraints( std::move(tile_size_maps), std::move(root_infos), - std::move(custom_constraints), + std::move(transpose_infos), std::move(custom_constraints), /*root_shape=*/instructions.back()->hlo()->shape(), device_description, std::move(tiled_emitter_constraints)))); }; @@ -258,6 +283,27 @@ absl::StatusOr TritonEmitterConstraints::ParametersSatisfyConstraints( } } + // Verify that no transpose op would require more shared memory than the + // device provides. A transpose stages its (padded) operand tile in shared + // memory to perform the layout conversion, so estimate that usage and reject + // tiles that would exceed the device shared memory limit. This lets the tile + // search fall back to a smaller tile instead of failing later with a + // RESOURCE_EXHAUSTED error during Triton compilation. + const int64_t shared_memory_limit = + device_info_.shared_memory_per_block_optin(); + for (const auto& transpose : transposes_) { + llvm::SmallVector operand_tile_sizes = + transpose.operand_size_map.Evaluate(tile_parameters); + if (GetPaddedTileSizeInBytes(operand_tile_sizes, + transpose.element_byte_size) > + shared_memory_limit) { + VLOG(2) << "Found a transpose whose operand tile would exceed the shared " + "memory limit of " + << shared_memory_limit << " bytes. Bailing out."; + return false; + } + } + return tiled_emitter_constraints_->ParametersSatisfyConstraints( tile_parameters); } @@ -349,6 +395,44 @@ Decision VerifyTritonConstraints(const TiledHloComputation& tiled_computation, } } + // 3. Transpose Shared Memory Limit. + // A transpose op stages its (padded) operand tile in shared memory to perform + // the layout conversion (the classic shared-memory transpose that avoids + // uncoalesced global memory accesses). Estimate that shared-memory usage as + // the product of the power-of-2-padded operand tile sizes multiplied by the + // element byte size, and reject tiles that would exceed the device shared + // memory limit. This lets the tile search fall back to a smaller tile instead + // of failing later with a RESOURCE_EXHAUSTED error during Triton compilation. + const int64_t shared_memory_limit = + device_info.shared_memory_per_block_optin(); + for (const TiledHloInstruction* inst : all_instructions) { + if (inst->hlo()->opcode() != HloOpcode::kTranspose) { + continue; + } + // The transposed operand (operand 0) is the tile that is staged in shared + // memory. + CHECK_EQ(inst->operands().size(), 1) + << "Transpose " << inst->hlo()->name() << " should have exactly one " + << "operand, but has " << inst->operands().size() << "."; + const TiledHloInstruction* operand = inst->operands().front(); + auto operand_tile_sizes_or = operand->tile().GetStaticTileSizes(); + if (!operand_tile_sizes_or.ok()) { + return Decision(operand_tile_sizes_or.status()); + } + + const int64_t shared_memory_bytes = GetPaddedTileSizeInBytes( + *operand_tile_sizes_or, ShapeUtil::ByteSizeOfPrimitiveType( + operand->hlo()->shape().element_type())); + if (shared_memory_bytes > shared_memory_limit) { + return Decision::Forbid(absl::StrCat( + "Transpose instruction ", inst->hlo()->name(), + " has an operand tile that requires an estimated ", + shared_memory_bytes, + " bytes of shared memory, which exceeds the device limit of ", + shared_memory_limit, " bytes.")); + } + } + VLOG(2) << "VerifyTritonConstraints: checking roots. Count: " << tiled_computation.roots().size(); for (const TiledHloInstruction* root : tiled_computation.roots()) { diff --git a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.h b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.h index bdba4e32021aac..ff45c75d8fc03c 100644 --- a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.h +++ b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints.h @@ -66,14 +66,23 @@ class TritonEmitterConstraints : public EmitterSpecificConstraints { std::vector dim_sizes; }; + // Holds the info needed to estimate the shared memory required to stage the + // operand tile of a transpose instruction. + struct TransposeTileInfo { + SymbolicMap operand_size_map; + int64_t element_byte_size; + }; + explicit TritonEmitterConstraints( llvm::SmallVector tile_size_maps, llvm::SmallVector roots, + llvm::SmallVector transposes, std::vector custom_constraints, const Shape& root_shape, const se::DeviceDescription& device_info, std::unique_ptr tiled_emitter_constraints) : tile_size_maps_(std::move(tile_size_maps)), roots_(std::move(roots)), + transposes_(std::move(transposes)), custom_constraints_(std::move(custom_constraints)), root_shape_(root_shape), device_info_(device_info), @@ -112,6 +121,10 @@ class TritonEmitterConstraints : public EmitterSpecificConstraints { // sizes evaluate to powers of 2 or have the same size as the dimension. llvm::SmallVector roots_; + // Holds the info for all transpose instructions necessary to estimate the + // shared memory required to stage their operand tiles. + llvm::SmallVector transposes_; + // Custom emitter-specific constraints to check in // `ParametersSatisfyConstraints`. std::vector custom_constraints_; diff --git a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints_test.cc b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints_test.cc index fe67ee81ead6c9..f5aa6420d42826 100644 --- a/third_party/xla/xla/service/gpu/model/triton_emitter_constraints_test.cc +++ b/third_party/xla/xla/service/gpu/model/triton_emitter_constraints_test.cc @@ -177,6 +177,77 @@ ENTRY entry_computation { absl_testing::IsOkAndHolds(false)); } +TEST_F(TritonEmitterConstraintsTest, + TransposeSharedMemoryConstraintIsEnforced) { + ASSERT_OK_AND_ASSIGN(std::unique_ptr module, + ParseAndReturnVerifiedModule(R"( +HloModule m + +fused_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT transpose = f32[1024,1024] transpose(param_0), dimensions={1,0} +} + +ENTRY entry_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT fusion = f32[1024,1024] fusion(param_0), kind=kCustom, + calls=fused_computation, backend_config={"fusion_backend_config":{"kind":"__triton"}} +} +)")); + + std::optional analysis = TryAnalyzeModule(module.get()); + ASSERT_TRUE(analysis.has_value()); + const HloInstruction* fusion_root = + module->entry_computation()->root_instruction()->fused_expression_root(); + + // The RTX A6000 test device has 99 * 1024 = 101376 bytes of opt-in shared + // memory. A [128, 128] f32 transpose operand tile requires + // 128 * 128 * 4 = 65536 bytes, which fits. + EXPECT_THAT(analysis->ParametersSatisfyConstraints( + Tiling({{fusion_root, FlatTiling({128, 128})}})), + absl_testing::IsOkAndHolds(true)); + + // A [256, 128] f32 transpose operand tile requires + // 256 * 128 * 4 = 131072 bytes, which exceeds the shared memory limit. The + // tile is otherwise valid (32768 elements is below the tensor size limit, + // tile sizes are powers of 2, and the number of blocks fits on the grid), so + // it is the transpose shared memory constraint that rejects it. + EXPECT_THAT(analysis->ParametersSatisfyConstraints( + Tiling({{fusion_root, FlatTiling({256, 128})}})), + absl_testing::IsOkAndHolds(false)); +} + +TEST_F(TritonEmitterConstraintsTest, NonTransposeIsNotChargedToSharedMemory) { + ASSERT_OK_AND_ASSIGN(std::unique_ptr module, + ParseAndReturnVerifiedModule(R"( +HloModule m + +fused_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT log = f32[1024,1024] log(param_0) +} + +ENTRY entry_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT fusion = f32[1024,1024] fusion(param_0), kind=kCustom, + calls=fused_computation, backend_config={"fusion_backend_config":{"kind":"__triton"}} +} +)")); + + std::optional analysis = TryAnalyzeModule(module.get()); + ASSERT_TRUE(analysis.has_value()); + const HloInstruction* fusion_root = + module->entry_computation()->root_instruction()->fused_expression_root(); + + // A [256, 128] f32 tile requires 256 * 128 * 4 = 131072 bytes, which would + // exceed the shared memory limit if it were charged to shared memory. Because + // there is no transpose op, the shared memory constraint does not apply and + // the tiling is accepted. + EXPECT_THAT(analysis->ParametersSatisfyConstraints( + Tiling({{fusion_root, FlatTiling({256, 128})}})), + absl_testing::IsOkAndHolds(true)); +} + TEST_F(TritonEmitterConstraintsTest, TooManyBlocksConstraintIsEnforced) { ASSERT_OK_AND_ASSIGN(std::unique_ptr module, ParseAndReturnVerifiedModule(R"( @@ -428,6 +499,63 @@ ENTRY entry_computation { StatusIs(_, HasSubstr("exceeds the maximum MMA dimension size"))); } +TEST_F(VerifyTritonConstraintsTest, TransposeSharedMemoryConstraintIsEnforced) { + ASSERT_OK_AND_ASSIGN(std::unique_ptr module, + ParseAndReturnVerifiedModule(R"hlo( +HloModule m + +fused_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT transpose = f32[1024,1024] transpose(param_0), dimensions={1,0} +} + +ENTRY entry_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT fusion = f32[1024,1024] fusion(param_0), kind=kCustom, + calls=fused_computation, + backend_config={"fusion_backend_config":{"kind":"__triton"}} +} +)hlo")); + + // The RTX A6000 test device has 99 * 1024 = 101376 bytes of opt-in shared + // memory. A [128, 128] f32 transpose operand tile requires + // 128 * 128 * 4 = 65536 bytes, which fits. + EXPECT_OK(CheckTiling(module.get(), {128, 128})); + + // A [256, 128] f32 transpose operand tile requires + // 256 * 128 * 4 = 131072 bytes, which exceeds the shared memory limit. The + // tile is otherwise valid (32768 elements is below the tensor size limit, + // tile sizes are powers of 2, and the number of blocks fits on the grid), so + // it is the transpose shared memory constraint that rejects it. + EXPECT_THAT(CheckTiling(module.get(), {256, 128}), + StatusIs(_, HasSubstr("exceeds the device limit"))); +} + +TEST_F(VerifyTritonConstraintsTest, NonTransposeIsNotChargedToSharedMemory) { + ASSERT_OK_AND_ASSIGN(std::unique_ptr module, + ParseAndReturnVerifiedModule(R"hlo( +HloModule m + +fused_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT log = f32[1024,1024] log(param_0) +} + +ENTRY entry_computation { + param_0 = f32[1024,1024] parameter(0) + ROOT fusion = f32[1024,1024] fusion(param_0), kind=kCustom, + calls=fused_computation, + backend_config={"fusion_backend_config":{"kind":"__triton"}} +} +)hlo")); + + // A [256, 128] f32 tile requires 256 * 128 * 4 = 131072 bytes, which would + // exceed the shared memory limit if it were charged to shared memory. Because + // there is no transpose op, the shared memory constraint does not apply and + // the tiling is accepted. + EXPECT_OK(CheckTiling(module.get(), {256, 128})); +} + TEST_F(VerifyTritonConstraintsTest, TooManyBlocksConstraintIsEnforced) { ASSERT_OK_AND_ASSIGN(std::unique_ptr module, ParseAndReturnVerifiedModule(R"hlo( From be9bceebc2958979dba070283089d12c3789d932 Mon Sep 17 00:00:00 2001 From: Parker Schuh Date: Tue, 1 Sep 2026 13:18:12 -0700 Subject: [PATCH 36/43] Allow using higher alignment for dmas for dynamic shape metadata. PiperOrigin-RevId: 974670040 --- third_party/xla/xla/pjrt/common_pjrt_client.cc | 7 +++++-- third_party/xla/xla/pjrt/dynamic_shapes.cc | 9 +++++++-- third_party/xla/xla/pjrt/dynamic_shapes.h | 3 ++- third_party/xla/xla/pjrt/raw_pjrt_client.h | 3 +++ 4 files changed, 17 insertions(+), 5 deletions(-) diff --git a/third_party/xla/xla/pjrt/common_pjrt_client.cc b/third_party/xla/xla/pjrt/common_pjrt_client.cc index eca3595ad77e00..6a63332acaf787 100644 --- a/third_party/xla/xla/pjrt/common_pjrt_client.cc +++ b/third_party/xla/xla/pjrt/common_pjrt_client.cc @@ -3867,12 +3867,15 @@ absl::StatusOr CommonPjRtBufferImpl::logical_on_device_shape() { PjRtDeviceEventRefVector definition_events) -> absl::StatusOr { auto ds_kind = client()->GetDynamicShapeKind(memory_space()->kind_id()); + size_t host_alignment_bytes = + buf_client->raw_client()->GetDmaHostAlignment(); PjRtDeviceEventSpan deps_span(definition_events); xla::ExecuteWhenReady( deps_span, buf_client->async_work_runner(), [definition_events = std::move(definition_events), raw_buffer = raw_buffer, output_shape = output_shape, - device_shape = std::move(device_shape), ds_kind]() mutable { + device_shape = std::move(device_shape), ds_kind, + host_alignment_bytes]() mutable { tsl::profiler::TraceMe traceme("D2H Read Shape Metadata"); absl::Status status = xla::GetErrors(definition_events); if (!status.ok()) { @@ -3883,7 +3886,7 @@ absl::StatusOr CommonPjRtBufferImpl::logical_on_device_shape() { return; } xla::ReadDynamicShape(raw_buffer, output_shape, device_shape, - ds_kind); + ds_kind, host_alignment_bytes); }); tsl::BlockUntilReady(output_shape.CopyRCRef().get()); if (auto* error = output_shape.GetErrorIfPresent()) { diff --git a/third_party/xla/xla/pjrt/dynamic_shapes.cc b/third_party/xla/xla/pjrt/dynamic_shapes.cc index fc39a352a67ef0..a81e7d60574c67 100644 --- a/third_party/xla/xla/pjrt/dynamic_shapes.cc +++ b/third_party/xla/xla/pjrt/dynamic_shapes.cc @@ -15,7 +15,10 @@ limitations under the License. #include "xla/pjrt/dynamic_shapes.h" +#include +#include #include +#include #include #include "absl/status/status.h" @@ -155,7 +158,8 @@ absl::StatusOr RemoveDynamicShapeMetadataIfPresent( void ReadDynamicShape(PjRtRawBufferRef raw_buffer, tsl::AsyncValueRef output_shape, - xla::Shape shape, PjRtDynamicShapeKind kind) { + xla::Shape shape, PjRtDynamicShapeKind kind, + size_t host_alignment_bytes) { auto requirements = PjRtShapeAndMetadataTransferRequirements::Get(shape, kind); if (requirements.metadata_size == 0) { @@ -180,7 +184,8 @@ void ReadDynamicShape(PjRtRawBufferRef raw_buffer, void* scratch = tsl::port::AlignedMalloc( requirements.metadata_size, - static_cast(requirements.metadata_alignment)); + static_cast( + std::max(requirements.metadata_alignment, host_alignment_bytes))); if (scratch == nullptr) { output_shape.SetError(absl::ResourceExhaustedError("AlignedMalloc failed")); return; diff --git a/third_party/xla/xla/pjrt/dynamic_shapes.h b/third_party/xla/xla/pjrt/dynamic_shapes.h index 2c36a774c94521..e8061dd7c0d1b1 100644 --- a/third_party/xla/xla/pjrt/dynamic_shapes.h +++ b/third_party/xla/xla/pjrt/dynamic_shapes.h @@ -75,7 +75,8 @@ absl::StatusOr RemoveDynamicShapeMetadataIfPresent( // Reads dynamic shape metadata into an output AsyncValueRef. void ReadDynamicShape(PjRtRawBufferRef raw_buffer, tsl::AsyncValueRef output_shape, - xla::Shape shape, PjRtDynamicShapeKind kind); + xla::Shape shape, PjRtDynamicShapeKind kind, + size_t host_alignment_bytes = 1); // Strips any metadata to give a logical shape. void StripMetadataForLogicalShape(xla::Shape& shape); diff --git a/third_party/xla/xla/pjrt/raw_pjrt_client.h b/third_party/xla/xla/pjrt/raw_pjrt_client.h index ca4cf11cef9dd8..0bce5ae8f8a7ef 100644 --- a/third_party/xla/xla/pjrt/raw_pjrt_client.h +++ b/third_party/xla/xla/pjrt/raw_pjrt_client.h @@ -183,6 +183,9 @@ class PjRtRawClient { return absl::UnimplementedError("DmaUnmap is not supported."); } + // Returns the required byte alignment for host memory when performing DMA. + virtual size_t GetDmaHostAlignment() const { return 1; } + virtual void UpdateGlobalProcessInfo( absl::Span infos) { LOG(WARNING) << "UpdateGlobalProcessInfo is not supported."; From 1e6ebb94a1afd392782472e8934fd2fefcc1adb9 Mon Sep 17 00:00:00 2001 From: Theotime Combes Date: Tue, 1 Sep 2026 13:23:09 -0700 Subject: [PATCH 37/43] [XLA] Add xla_deduplicate_backend_configs_min_size flag to control backend config deduplication Expose `xla_deduplicate_backend_configs_min_size` in `DebugOptions`, defaulting to `MAX_INT` (disabled). Setting this flag to any valid non-negative threshold enables backend config deduplication into payloads at serialization time in `HloModule::ToProto()` across TPU, GPU, CPU, etc., without requiring `HloProtoOptions` to be threaded through downstream compiler APIs. Callers that manually enable deduplication (such as `gpu_executable.cc`) are preserved. PiperOrigin-RevId: 974672733 --- third_party/xla/xla/debug_options_flags.cc | 10 ++++ .../xla/xla/debug_options_flags_test.cc | 18 ++++++ third_party/xla/xla/hlo/ir/hlo_module.cc | 13 +++++ third_party/xla/xla/hlo/ir/hlo_module_test.cc | 58 +++++++++++++++++++ third_party/xla/xla/xla.proto | 7 ++- 5 files changed, 105 insertions(+), 1 deletion(-) diff --git a/third_party/xla/xla/debug_options_flags.cc b/third_party/xla/xla/debug_options_flags.cc index 3c0289ec0c7ea0..8d36c0bf83f141 100644 --- a/third_party/xla/xla/debug_options_flags.cc +++ b/third_party/xla/xla/debug_options_flags.cc @@ -483,6 +483,8 @@ DebugOptions DefaultDebugOptionsIgnoringFlags() { opts.set_xla_gpu_use_new_autotune_cache_format(true); opts.set_xla_compile_all_supported_configs(false); + opts.set_xla_deduplicate_backend_configs_min_size( + std::numeric_limits::max()); opts.set_xla_gpu_experimental_autotune_cache_mode( DebugOptions::AUTOTUNE_CACHE_MODE_UPDATE); @@ -3050,6 +3052,14 @@ void MakeDebugOptionsFlags(std::vector* flag_list, debug_options->xla_compile_all_supported_configs(), "When autotuning is disabled, if true, compiles all supported configs" " in parallel before returning the first successful one.")); + flag_list->push_back(tsl::Flag( + "xla_deduplicate_backend_configs_min_size", + int64_setter_for( + &DebugOptions::set_xla_deduplicate_backend_configs_min_size), + debug_options->xla_deduplicate_backend_configs_min_size(), + "Minimum backend_config size (in bytes) to be eligible for deduplication " + "into payloads during serialization. Configs smaller than this threshold " + "are kept inline. Default is MAX_INT (feature disabled).")); flag_list->push_back(tsl::Flag( "xla_gpu_experimental_autotune_backends", SetterForRepeatedEnum( diff --git a/third_party/xla/xla/debug_options_flags_test.cc b/third_party/xla/xla/debug_options_flags_test.cc index ad8a6ebf838a87..5d9f3ddbb58127 100644 --- a/third_party/xla/xla/debug_options_flags_test.cc +++ b/third_party/xla/xla/debug_options_flags_test.cc @@ -15,6 +15,8 @@ limitations under the License. #include "xla/debug_options_flags.h" +#include +#include #include #include #include @@ -462,5 +464,21 @@ TEST(DebugOptions, DisableHloPassesRejectsMalformedEntries) { EXPECT_FALSE(ParseEnableHloPassesOnlyFlag("@").first); } +TEST(DebugOptions, DeduplicateBackendConfigsMinSizeDefaultIsMaxInt) { + DebugOptions opts = DefaultDebugOptionsIgnoringFlags(); + EXPECT_EQ(opts.xla_deduplicate_backend_configs_min_size(), + std::numeric_limits::max()); +} + +TEST(DebugOptions, DeduplicateBackendConfigsMinSizeFlagsParsing) { + DebugOptions opts; + std::vector flags; + MakeDebugOptionsFlags(&flags, &opts); + std::vector flag_args = { + "--xla_deduplicate_backend_configs_min_size=128"}; + EXPECT_TRUE(tsl::Flags::Parse(flag_args, flags)); + EXPECT_EQ(opts.xla_deduplicate_backend_configs_min_size(), 128); +} + } // namespace } // namespace xla diff --git a/third_party/xla/xla/hlo/ir/hlo_module.cc b/third_party/xla/xla/hlo/ir/hlo_module.cc index c8fa7c566d2db2..179e1d8d41d0ab 100644 --- a/third_party/xla/xla/hlo/ir/hlo_module.cc +++ b/third_party/xla/xla/hlo/ir/hlo_module.cc @@ -21,6 +21,7 @@ limitations under the License. #include #include #include +#include #include #include #include @@ -615,6 +616,18 @@ void HloModule::ToProto(HloModuleProto* proto, HloProtoOptions options) const { entry_computation_layout().ComputeProgramShape().ToProto(); } + // Deduplicate backend configs if requested via options or via XLA flag. + // Do not override options where it's already manually set by the caller. + if (!options.deduplicate_backend_config && + config().debug_options().has_xla_deduplicate_backend_configs_min_size()) { + int64_t min_size = + config().debug_options().xla_deduplicate_backend_configs_min_size(); + if (min_size >= 0 && min_size < std::numeric_limits::max()) { + options.deduplicate_backend_config = true; + options.min_backend_config_size = min_size; + } + } + // Instantiate one shared deduplicator when either option is enabled. std::optional payload_deduplicator; if (options.deduplicate_backend_config || options.deduplicate_metadata) { diff --git a/third_party/xla/xla/hlo/ir/hlo_module_test.cc b/third_party/xla/xla/hlo/ir/hlo_module_test.cc index e747795bfd523a..e728342a75932b 100644 --- a/third_party/xla/xla/hlo/ir/hlo_module_test.cc +++ b/third_party/xla/xla/hlo/ir/hlo_module_test.cc @@ -2059,6 +2059,64 @@ TEST(HloModuleTest, BackendConfigDeduplicationRespectsMinSize) { EXPECT_EQ(large->backend_config_payload().id(), 0); } +TEST(HloModuleTest, BackendConfigDeduplicationViaDebugOptionsFlag) { + const char* hlo_text = R"( + HloModule test_module + ENTRY comp { + p0 = f32[] parameter(0) + p1 = f32[] parameter(1) + ROOT add = f32[] add(p0, p1) + })"; + ASSERT_OK_AND_ASSIGN(auto m, ParseAndReturnUnverifiedModule(hlo_text)); + HloInstruction* p0 = m->entry_computation()->GetInstructionWithName("p0"); + HloInstruction* p1 = m->entry_computation()->GetInstructionWithName("p1"); + + std::string small_config = "short"; // 5 bytes + std::string large_config = std::string(200, 'x'); // 200 bytes + + p0->set_raw_backend_config_string(small_config); + p1->set_raw_backend_config_string(large_config); + + // By default (min_size == MAX_INT), ToProto() does NOT deduplicate configs. + HloModuleProto default_proto = m->ToProto(); + EXPECT_EQ(default_proto.payloads_size(), 0); + + // Enable via DebugOptions flag with min size threshold = 128. + m->mutable_config() + .mutable_debug_options() + .set_xla_deduplicate_backend_configs_min_size(128); + + // Default ToProto() should now deduplicate large_config. + HloModuleProto proto = m->ToProto(); + + ASSERT_EQ(proto.payloads_size(), 1); + EXPECT_EQ(proto.payloads(0), large_config); + + const auto& instructions = proto.computations(0).instructions(); + const auto* small = &instructions[0]; + const auto* large = &instructions[1]; + if (small->name() != "p0") { + std::swap(small, large); + } + + // Small config stays inline. + EXPECT_EQ(small->backend_config(), small_config); + EXPECT_FALSE(small->has_backend_config_payload()); + + // Large config is deduplicated into a payload. + EXPECT_EQ(large->backend_config(), ""); + EXPECT_TRUE(large->has_backend_config_payload()); + EXPECT_EQ(large->backend_config_payload().id(), 0); + + // Verify that explicit HloProtoOptions overrides are preserved: + // e.g. caller explicitly sets min_backend_config_size to 256, so large_config + // (200 bytes) is kept inline. + HloModuleProto proto_manual = m->ToProto(HloProtoOptions{ + /*deduplicate_backend_config=*/true, /*deduplicate_metadata=*/false, + /*min_backend_config_size=*/256}); + EXPECT_EQ(proto_manual.payloads_size(), 0); +} + TEST(HloModuleTest, BackendConfigNoInternByDefault) { const char* hlo_text = R"( HloModule test_module diff --git a/third_party/xla/xla/xla.proto b/third_party/xla/xla/xla.proto index b9d9e0374324f3..e8e40de0502825 100644 --- a/third_party/xla/xla/xla.proto +++ b/third_party/xla/xla/xla.proto @@ -1832,7 +1832,12 @@ message DebugOptions { repeated CollectiveOpType xla_gpu_unsupported_use_cross_host_one_shot_kernel = 541; - // Next id: 542 + // Minimum backend_config size (in bytes) to be eligible for deduplication + // into payloads during serialization. Configs smaller than this threshold are + // kept inline. Default is MAX_INT (deduplication disabled). + optional int64 xla_deduplicate_backend_configs_min_size = 542; + + // Next id: 543 // Extra options to pass to the compilation backend (e.g. LLVM); specific // interpretation of these values is left to the backend. From aa0530d5f5864b6f634fdfe21f7a329ea6b4ff12 Mon Sep 17 00:00:00 2001 From: Shyamli Agrawal Date: Tue, 1 Sep 2026 14:08:25 -0700 Subject: [PATCH 38/43] Remove the no-op flow of uploading and using autotune results. PiperOrigin-RevId: 974697845 --- third_party/xla/xla/service/gpu/BUILD | 1 - third_party/xla/xla/service/gpu/export_hlo.h | 7 +++---- third_party/xla/xla/service/gpu/gpu_compiler.cc | 3 +-- 3 files changed, 4 insertions(+), 7 deletions(-) diff --git a/third_party/xla/xla/service/gpu/BUILD b/third_party/xla/xla/service/gpu/BUILD index a15a040a367045..1f6a4aee2d3fce 100644 --- a/third_party/xla/xla/service/gpu/BUILD +++ b/third_party/xla/xla/service/gpu/BUILD @@ -3923,7 +3923,6 @@ cc_library( hdrs = ["export_hlo.h"], compatible_with = get_compatible_with_portable(), deps = [ - "//xla:autotune_results_proto_cc", "//xla/hlo/ir:hlo", "//xla/stream_executor:device_description_proto_cc", "@com_google_absl//absl/strings", diff --git a/third_party/xla/xla/service/gpu/export_hlo.h b/third_party/xla/xla/service/gpu/export_hlo.h index 2af0e31f91a54e..073d1dc60f6639 100644 --- a/third_party/xla/xla/service/gpu/export_hlo.h +++ b/third_party/xla/xla/service/gpu/export_hlo.h @@ -22,7 +22,6 @@ limitations under the License. #include #include "absl/strings/string_view.h" -#include "xla/autotune_results.pb.h" #include "xla/hlo/ir/hlo_module.h" #include "xla/stream_executor/device_description.pb.h" @@ -46,7 +45,7 @@ class SymbolUploader { const stream_executor::GpuTargetConfigProto& gpu_target_config) = 0; virtual std::optional MaybeUploadOptimizedHloModule( - HloModule* module, const AutotuneResults& autotune_results) = 0; + HloModule* module) = 0; virtual void MaybeUploadSymbolMapping( absl::string_view unoptimized_fingerprint, @@ -89,10 +88,10 @@ inline std::optional MaybeUploadUnoptimizedGpuSymbols( } inline std::optional MaybeUploadOptimizedGpuSymbols( - HloModule* module, const AutotuneResults& autotune_results) { + HloModule* module) { if (SymbolUploader* uploader = GetGlobalSymbolUploaderRegistry().uploader(); uploader != nullptr) { - return uploader->MaybeUploadOptimizedHloModule(module, autotune_results); + return uploader->MaybeUploadOptimizedHloModule(module); } return std::nullopt; diff --git a/third_party/xla/xla/service/gpu/gpu_compiler.cc b/third_party/xla/xla/service/gpu/gpu_compiler.cc index 3f12c4b13efb2d..a14075073a6126 100644 --- a/third_party/xla/xla/service/gpu/gpu_compiler.cc +++ b/third_party/xla/xla/service/gpu/gpu_compiler.cc @@ -2367,8 +2367,7 @@ absl::StatusOr> GpuCompiler::RunHloPasses( } std::optional optimized_fingerprint; if (should_upload_hlo_modules) { - optimized_fingerprint = - MaybeUploadOptimizedGpuSymbols(module.get(), AutotuneResults()); + optimized_fingerprint = MaybeUploadOptimizedGpuSymbols(module.get()); } if (unoptimized_fingerprint.has_value() && optimized_fingerprint.has_value()) { From f0f71dc7f83185b30e32cbe676cd0ec258f6c31b Mon Sep 17 00:00:00 2001 From: Zac Mustin Date: Tue, 1 Sep 2026 14:33:16 -0700 Subject: [PATCH 39/43] Replace `TF_ASSERT_OK_AND_ASSIGN` with `ASSERT_OK_AND_ASSIGN` in `pjrt_c_api_client_test.cc`. `TF_ASSERT_OK_AND_ASSIGN` has been marked deprecated. PiperOrigin-RevId: 974711143 --- third_party/xla/xla/pjrt/c_api_client/BUILD | 1 - .../c_api_client/pjrt_c_api_client_test.cc | 342 +++++++++--------- 2 files changed, 167 insertions(+), 176 deletions(-) diff --git a/third_party/xla/xla/pjrt/c_api_client/BUILD b/third_party/xla/xla/pjrt/c_api_client/BUILD index cefc24a429f38c..9005ed67b09eb1 100644 --- a/third_party/xla/xla/pjrt/c_api_client/BUILD +++ b/third_party/xla/xla/pjrt/c_api_client/BUILD @@ -183,7 +183,6 @@ xla_cc_test( "//xla/tsl/lib/core:status_test_util", "//xla/tsl/platform:env", "//xla/tsl/platform:errors", - "//xla/tsl/platform:statusor", "//xla/tsl/platform:test", "@com_google_absl//absl/status", "@com_google_absl//absl/status:status_matchers", diff --git a/third_party/xla/xla/pjrt/c_api_client/pjrt_c_api_client_test.cc b/third_party/xla/xla/pjrt/c_api_client/pjrt_c_api_client_test.cc index ade4ee76e0808c..3c47f0ea8e8a07 100644 --- a/third_party/xla/xla/pjrt/c_api_client/pjrt_c_api_client_test.cc +++ b/third_party/xla/xla/pjrt/c_api_client/pjrt_c_api_client_test.cc @@ -64,7 +64,6 @@ limitations under the License. #include "xla/tsl/lib/core/status_test_util.h" #include "xla/tsl/platform/env.h" #include "xla/tsl/platform/errors.h" -#include "xla/tsl/platform/statusor.h" #include "xla/tsl/platform/test.h" #include "xla/tsl/platform/threadpool.h" #include "xla/types.h" @@ -90,17 +89,17 @@ static void SetUpCpuPjRtApi() { TEST(PjRtCApiClientTest, FulfillAliasBuffer) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); std::vector data{1, 2, 3, 4, 5, 6}; Shape shape = ShapeUtil::MakeShape(S32, {2, 3}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto alias_buffer, client->CreateAliasBuffer(shape, client->memory_spaces()[0])); // Create a buffer from host data. - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto param, client->BufferFromHostBuffer( data.data(), shape.element_type(), shape.dimensions(), @@ -116,12 +115,11 @@ TEST(PjRtCApiClientTest, FulfillAliasBuffer) { auto computation = builder.Build(add).value(); // Compile and load the executable. - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr executable, - client->CompileAndLoad(computation, CompileOptions())); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, CompileOptions())); // Execute the kernel. - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector>> results, executable->Execute({{param.get()}}, ExecuteOptions())); ASSERT_EQ(results.size(), 1); @@ -134,8 +132,8 @@ TEST(PjRtCApiClientTest, FulfillAliasBuffer) { // Fulfill the alias buffer with the result of the add one kernel. ASSERT_NE(alias_buffer.second, nullptr); TF_ASSERT_OK(std::move(alias_buffer.second)(result_buffer.get())); - TF_ASSERT_OK_AND_ASSIGN(auto alias_literal, - alias_buffer.first->ToLiteral().Await()); + ASSERT_OK_AND_ASSIGN(auto alias_literal, + alias_buffer.first->ToLiteral().Await()); // Expected result: data + 1 EXPECT_TRUE(LiteralTestUtil::Equal( @@ -144,14 +142,14 @@ TEST(PjRtCApiClientTest, FulfillAliasBuffer) { TEST(PjRtCApiClientTest, CreateErrorBuffer) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); absl::Status error = absl::InternalError("Test Error"); error.SetPayload("test_key", absl::Cord("test_payload_value")); Shape shape = ShapeUtil::MakeShape(S32, {2, 3}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto error_buffer, client->CreateErrorBuffer(error, shape, client->memory_spaces()[0])); @@ -164,8 +162,8 @@ TEST(PjRtCApiClientTest, CreateErrorBuffer) { TEST(PjRtCApiClientTest, ConcurrentGetReadyFuture) { const PJRT_Api* c_api = ::pjrt::cpu_plugin::GetCpuPjrtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - WrapClientAroundCApi(c_api)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + WrapClientAroundCApi(c_api)); constexpr int kNumThreads = 4; tsl::thread::ThreadPool thread_pool( @@ -175,7 +173,7 @@ TEST(PjRtCApiClientTest, ConcurrentGetReadyFuture) { Shape shape = ShapeUtil::MakeShape(S32, {2, 3}); // Create a buffer from host data. - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto param, client->BufferFromHostBuffer( data.data(), shape.element_type(), shape.dimensions(), @@ -191,11 +189,10 @@ TEST(PjRtCApiClientTest, ConcurrentGetReadyFuture) { auto computation = builder.Build(add).value(); // Compile and load the executable. - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr executable, - client->CompileAndLoad(computation, CompileOptions())); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, CompileOptions())); for (size_t i = 0; i < 100; ++i) { - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector>> results, executable->Execute({{param.get()}}, ExecuteOptions())); auto buffer = std::move(results[0][0]); @@ -213,13 +210,13 @@ TEST(PjRtCApiClientTest, ConcurrentGetReadyFuture) { TEST(PjRtCApiClientTest, GetReadyFutureDeletedBuffer) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); std::vector data{1}; Shape shape = ShapeUtil::MakeShape(S32, {}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::unique_ptr buffer, client->BufferFromHostBuffer( data.data(), shape.element_type(), shape.dimensions(), @@ -237,12 +234,12 @@ TEST(PjRtCApiClientTest, GetReadyFutureDeletedBuffer) { TEST(PjRtCApiClientTest, IsDynamicDimension) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); // Prepare input buffer and executable. std::vector data0{1, 2, 3, 4, 5, 6}; Shape shape0 = ShapeUtil::MakeShape(S32, {2, 3}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto param0, client->BufferFromHostBuffer( data0.data(), shape0.element_type(), shape0.dimensions(), @@ -251,7 +248,7 @@ TEST(PjRtCApiClientTest, IsDynamicDimension) { client->memory_spaces()[0], /*device_layout=*/nullptr)); std::vector data1{2}; Shape shape1 = ShapeUtil::MakeShape(S32, {}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto param1, client->BufferFromHostBuffer( data1.data(), shape1.element_type(), shape1.dimensions(), @@ -350,12 +347,12 @@ TEST(PjRtCApiClientTest, DynamicShapesPipeline) { TEST(PjRtCApiClientTest, OnDeviceShape) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); std::vector data{1, 2, 3, 4, 5, 6}; for (PrimitiveType t : {F32, F16, S8, BF16}) { Shape shape = ShapeUtil::MakeShape(t, {3, 2}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto buffer, client->BufferFromHostBuffer( data.data(), shape.element_type(), shape.dimensions(), @@ -369,8 +366,8 @@ TEST(PjRtCApiClientTest, OnDeviceShape) { TEST(PjRtCApiClientTest, ClientPlatformIdAndName) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); EXPECT_EQ(client->platform_name(), xla::CpuName()); EXPECT_EQ(client->platform_id(), xla::CpuId()); @@ -378,11 +375,11 @@ TEST(PjRtCApiClientTest, ClientPlatformIdAndName) { TEST(PjRtCApiClientTest, TopologyPlatformIdAndName) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); - TF_ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, - client->GetTopologyDescription()); + ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, + client->GetTopologyDescription()); ASSERT_NE(topology, nullptr); EXPECT_EQ(topology->platform_name(), xla::CpuName()); @@ -391,16 +388,16 @@ TEST(PjRtCApiClientTest, TopologyPlatformIdAndName) { TEST(PjRtCApiClientTest, TopologyGetDefaultLayout) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); - TF_ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, - client->GetTopologyDescription()); + ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, + client->GetTopologyDescription()); ASSERT_NE(topology, nullptr); std::vector dims = {2, 3, 4}; - TF_ASSERT_OK_AND_ASSIGN(Layout layout, - topology->GetDefaultLayout(PrimitiveType::F32, dims)); + ASSERT_OK_AND_ASSIGN(Layout layout, + topology->GetDefaultLayout(PrimitiveType::F32, dims)); Layout expected_layout = LayoutUtil::MakeDescendingLayout(dims.size()); EXPECT_EQ(layout, expected_layout); @@ -408,21 +405,21 @@ TEST(PjRtCApiClientTest, TopologyGetDefaultLayout) { TEST(PjRtCApiClientTest, TopologyFingerprint) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); - TF_ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, - client->GetTopologyDescription()); + ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, + client->GetTopologyDescription()); ASSERT_NE(topology, nullptr); - TF_ASSERT_OK_AND_ASSIGN(uint64_t fingerprint, topology->Fingerprint()); + ASSERT_OK_AND_ASSIGN(uint64_t fingerprint, topology->Fingerprint()); EXPECT_NE(fingerprint, 0); } TEST(PjRtCApiClientTest, NonEmptyExecutableFingerprint) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); Shape shape = ShapeUtil::MakeShapeWithType({4}); XlaBuilder builder("sum"); auto inp_0 = Parameter(&builder, 0, shape, "input0"); @@ -471,8 +468,8 @@ TEST(PjRtCApiClientTest, GetCompileOptions) { TEST(PjRtCApiClientTest, CreateBuffersForAsyncHostToDeviceWithShape) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); xla::Shape host_shape = xla::ShapeUtil::MakeShapeWithDenseLayout( xla::PrimitiveType::F32, /*dimensions=*/{2, 2, 2}, /*minor_to_major=*/{1, 0, 2}); @@ -485,15 +482,15 @@ TEST(PjRtCApiClientTest, CreateBuffersForAsyncHostToDeviceWithShape) { TEST(PjRtClientTest, CreateViewAndCopyToDeviceAsyncExternalCpuOnly) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); ASSERT_GT(client->addressable_devices().size(), 1); alignas(cpu::MinAlign()) std::array data; data.fill(0); auto* data_ptr = data.data(); Shape shape = ShapeUtil::MakeShape(S32, {4}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto buffer, client->CreateViewOfDeviceBuffer( data_ptr, shape, client->memory_spaces()[0], @@ -501,12 +498,11 @@ TEST(PjRtClientTest, CreateViewAndCopyToDeviceAsyncExternalCpuOnly) { (void)data; })); - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr result, - buffer->CopyToMemorySpace(client->memory_spaces()[1])); + ASSERT_OK_AND_ASSIGN(std::unique_ptr result, + buffer->CopyToMemorySpace(client->memory_spaces()[1])); buffer.reset(); ASSERT_TRUE(result); - TF_ASSERT_OK_AND_ASSIGN(auto literal, result->ToLiteral().Await()); + ASSERT_OK_AND_ASSIGN(auto literal, result->ToLiteral().Await()); std::vector expected(4, 0); EXPECT_TRUE(LiteralTestUtil::Equal(LiteralUtil::CreateR1(expected), @@ -515,14 +511,14 @@ TEST(PjRtClientTest, CreateViewAndCopyToDeviceAsyncExternalCpuOnly) { TEST(PjRtClientTest, CompileUsesStableHloVersion) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(const PJRT_Api* c_api, pjrt::PjrtApi("cpu")); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(const PJRT_Api* c_api, pjrt::PjrtApi("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); static auto PJRT_Client_Compile_Orig = c_api->PJRT_Client_Compile; constexpr char kProgram[] = "func.func @main() {return}"; auto context = std::make_unique(); - TF_ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, - ParseMlirModuleString(kProgram, *context)); + ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, + ParseMlirModuleString(kProgram, *context)); const_cast(c_api)->PJRT_Client_Compile = [](PJRT_Client_Compile_Args* args) -> PJRT_Error* { mlir::vhlo::Version version = mlir::vhlo::Version::getCurrentVersion(); @@ -544,12 +540,12 @@ TEST(PjRtClientTest, CompileUsesStableHloVersion) { TEST(PjRtClientTest, CompileWorksInplace) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); constexpr char kProgram[] = "func.func @main() {return}"; auto context = std::make_unique(); - TF_ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, - ParseMlirModuleString(kProgram, *context)); + ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, + ParseMlirModuleString(kProgram, *context)); CompileOptions options; options.allow_in_place_mlir_modification = true; @@ -563,43 +559,40 @@ TEST(PjRtClientTest, CompileWorksInplace) { TEST(PjRtClientTest, CompileMlirModule) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); constexpr char kProgram[] = "func.func @main() {return}"; auto context = std::make_unique(); - TF_ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, - ParseMlirModuleString(kProgram, *context)); + ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, + ParseMlirModuleString(kProgram, *context)); CompileOptions options; - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr executable, - client->Compile( - MaybeOwningMlirModule(std::move(context), std::move(module)), - options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->Compile(MaybeOwningMlirModule(std::move(context), + std::move(module)), + options)); EXPECT_NE(executable.get(), nullptr); } TEST(PjRtCApiClientTest, LoadExecutable) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); constexpr char kProgram[] = "func.func @main() {return}"; auto context = std::make_unique(); - TF_ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, - ParseMlirModuleString(kProgram, *context)); + ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, + ParseMlirModuleString(kProgram, *context)); CompileOptions options; - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr executable, - client->Compile( - MaybeOwningMlirModule(std::move(context), std::move(module)), - options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->Compile(MaybeOwningMlirModule(std::move(context), + std::move(module)), + options)); ASSERT_NE(executable.get(), nullptr); - TF_ASSERT_OK_AND_ASSIGN( - std::unique_ptr loaded_executable, - client->Load(std::move(executable), LoadOptions{})); + ASSERT_OK_AND_ASSIGN(std::unique_ptr loaded_executable, + client->Load(std::move(executable), LoadOptions{})); ASSERT_NE(loaded_executable.get(), nullptr); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector>> results, loaded_executable->Execute(/*argument_handles=*/{{}}, ExecuteOptions())); ASSERT_EQ(results.size(), 1); @@ -608,28 +601,27 @@ TEST(PjRtCApiClientTest, LoadExecutable) { TEST(PjRtCApiClientTest, LoadSameExecutableTwice) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); constexpr char kProgram[] = "func.func @main() {return}"; auto context = std::make_unique(); - TF_ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, - ParseMlirModuleString(kProgram, *context)); + ASSERT_OK_AND_ASSIGN(mlir::OwningOpRef module, + ParseMlirModuleString(kProgram, *context)); CompileOptions options; - TF_ASSERT_OK_AND_ASSIGN( - const std::shared_ptr executable, - client->Compile( - MaybeOwningMlirModule(std::move(context), std::move(module)), - options)); + ASSERT_OK_AND_ASSIGN(const std::shared_ptr executable, + client->Compile(MaybeOwningMlirModule(std::move(context), + std::move(module)), + options)); ASSERT_NE(executable.get(), nullptr); // Load the executable twice. { - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::unique_ptr loaded_executable, client->Load(executable, LoadOptions{})); ASSERT_NE(loaded_executable.get(), nullptr); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector>> results, loaded_executable->Execute(/*argument_handles=*/{{}}, ExecuteOptions())); @@ -637,12 +629,12 @@ TEST(PjRtCApiClientTest, LoadSameExecutableTwice) { EXPECT_EQ(results[0].size(), 0); } { - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::unique_ptr loaded_executable, client->Load(executable, LoadOptions{})); ASSERT_NE(loaded_executable.get(), nullptr); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector>> results, loaded_executable->Execute(/*argument_handles=*/{{}}, ExecuteOptions())); @@ -653,10 +645,10 @@ TEST(PjRtCApiClientTest, LoadSameExecutableTwice) { TEST(PjRtClientTest, CanQueryMemoryDescriptions) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); - TF_ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, - client->GetTopologyDescription()); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, + client->GetTopologyDescription()); std::vector> devices = topology->DeviceDescriptions(); for (std::unique_ptr& device : devices) { @@ -672,8 +664,8 @@ TEST(PjRtClientTest, CanQueryMemoryDescriptions) { TEST(PjRtCApiClientTest, GetDeviceAssignment) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); ASSERT_GT(client->addressable_devices().size(), 1); XlaBuilder builder("Identity"); @@ -688,8 +680,8 @@ TEST(PjRtCApiClientTest, GetDeviceAssignment) { CompileOptions options; options.executable_build_options.set_device_assignment(device_assignment); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, - client->CompileAndLoad(computation, options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, options)); const DeviceAssignment& retrieved_assignment = executable->device_assignment(); @@ -701,8 +693,8 @@ TEST(PjRtCApiClientTest, GetDeviceAssignment) { TEST(PjRtCApiClientTest, WrapClientAroundCApi) { const PJRT_Api* c_api = ::pjrt::cpu_plugin::GetCpuPjrtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - WrapClientAroundCApi(c_api)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + WrapClientAroundCApi(c_api)); EXPECT_EQ(client->platform_name(), xla::CpuName()); EXPECT_EQ(client->platform_id(), xla::CpuId()); } @@ -741,12 +733,12 @@ TEST(PjRtCApiClientTest, ForwardExecuteContext) { })"; const PJRT_Api* c_api = ::pjrt::cpu_plugin::GetCpuPjrtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - WrapClientAroundCApi(c_api)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + WrapClientAroundCApi(c_api)); - TF_ASSERT_OK_AND_ASSIGN(auto hlo_module, - ParseAndReturnUnverifiedModule(kProgram, {})); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN(auto hlo_module, + ParseAndReturnUnverifiedModule(kProgram, {})); + ASSERT_OK_AND_ASSIGN( auto executable, client->CompileAndLoad(XlaComputation(hlo_module->ToProto()), {})); @@ -756,10 +748,11 @@ TEST(PjRtCApiClientTest, ForwardExecuteContext) { ExecuteOptions options; options.context = &context; - auto result = executable->Execute(/*argument_handles=*/{{}}, options); + ASSERT_OK_AND_ASSIGN(auto result, + executable->Execute(/*argument_handles=*/{{}}, options)); - TF_ASSERT_OK_AND_ASSIGN(std::shared_ptr result_literal, - result->at(0).at(0)->ToLiteral().Await()); + ASSERT_OK_AND_ASSIGN(std::shared_ptr result_literal, + result.at(0).at(0)->ToLiteral().Await()); EXPECT_TRUE(LiteralTestUtil::Equal( LiteralUtil::CreateR1({42.0f, 42.0f, 42.0f, 42.0f}), *result_literal)); @@ -767,8 +760,8 @@ TEST(PjRtCApiClientTest, ForwardExecuteContext) { TEST(PjRtClientTest, DeserializeExecutableWithDifferentDeviceAssignment) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); ASSERT_GT(client->addressable_devices().size(), 1); XlaBuilder builder("Identity"); @@ -788,11 +781,11 @@ TEST(PjRtClientTest, DeserializeExecutableWithDifferentDeviceAssignment) { std::unique_ptr executable = client->CompileAndLoad(computation, compile_options_for_device(0)) .value(); - TF_ASSERT_OK_AND_ASSIGN(std::string serialized_executable, - executable->SerializeExecutable()); + ASSERT_OK_AND_ASSIGN(std::string serialized_executable, + executable->SerializeExecutable()); // Deserialize the executable for device 1. - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto deserialized_executable, client->LoadSerializedExecutable( serialized_executable, compile_options_for_device(1), LoadOptions{})); @@ -813,17 +806,17 @@ TEST(PjRtCApiClientTest, GetOutputShapes) { })"; const PJRT_Api* c_api = ::pjrt::cpu_plugin::GetCpuPjrtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - WrapClientAroundCApi(c_api)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + WrapClientAroundCApi(c_api)); - TF_ASSERT_OK_AND_ASSIGN(auto hlo_module, - ParseAndReturnUnverifiedModule(kProgram, {})); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN(auto hlo_module, + ParseAndReturnUnverifiedModule(kProgram, {})); + ASSERT_OK_AND_ASSIGN( auto executable, client->CompileAndLoad(XlaComputation(hlo_module->ToProto()), {})); - TF_ASSERT_OK_AND_ASSIGN(std::vector output_shapes, - executable->GetOutputShapes()); + ASSERT_OK_AND_ASSIGN(std::vector output_shapes, + executable->GetOutputShapes()); EXPECT_EQ(output_shapes.size(), 1); Shape expected_shape = ShapeUtil::MakeShape(F32, {4}); EXPECT_EQ(output_shapes[0], expected_shape); @@ -831,8 +824,8 @@ TEST(PjRtCApiClientTest, GetOutputShapes) { TEST(PjRtCApiClientTest, GetParameterAndOutputShardings) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); Shape shape = ShapeUtil::MakeShapeWithType({4}); XlaBuilder builder("sum"); auto inp_0 = Parameter(&builder, 0, shape, "input0"); @@ -841,8 +834,8 @@ TEST(PjRtCApiClientTest, GetParameterAndOutputShardings) { auto computation = builder.Build(sum).value(); CompileOptions options; - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, - client->CompileAndLoad(computation, options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, options)); // CPU usually returns nullopt for shardings if not explicitly set. auto parameter_shardings = executable->GetParameterShardings(); @@ -860,8 +853,8 @@ TEST(PjRtCApiClientTest, GetParameterAndOutputShardings) { TEST(PjRtCApiClientTest, GetParameterAndOutputLayouts) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); Shape shape = ShapeUtil::MakeShapeWithType({4}); XlaBuilder builder("sum"); @@ -871,17 +864,17 @@ TEST(PjRtCApiClientTest, GetParameterAndOutputLayouts) { auto computation = builder.Build(sum).value(); CompileOptions options; - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, - client->CompileAndLoad(computation, options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, options)); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector> parameter_layouts, executable->GetParameterLayouts()); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::vector> output_layouts, executable->GetOutputLayouts()); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( xla::Layout expected_layout, client->GetDefaultLayout(shape.element_type(), shape.dimensions())); @@ -901,29 +894,28 @@ TEST(PjRtCApiClientTest, GetParameterAndOutputLayouts) { TEST(PjRtClientTest, BufferFromLiteralInt4) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); xla::Shape shape = xla::ShapeUtil::MakeShape(S4, {128, 256}); - TF_ASSERT_OK_AND_ASSIGN(auto literal, xla::MakeFakeLiteral(shape)); - TF_ASSERT_OK_AND_ASSIGN( - auto buffer, - client->BufferFromHostLiteral(literal, client->memory_spaces()[0])); - TF_ASSERT_OK_AND_ASSIGN(auto received_literal, buffer->ToLiteral().Await()); + ASSERT_OK_AND_ASSIGN(auto literal, xla::MakeFakeLiteral(shape)); + ASSERT_OK_AND_ASSIGN(auto buffer, client->BufferFromHostLiteral( + literal, client->memory_spaces()[0])); + ASSERT_OK_AND_ASSIGN(auto received_literal, buffer->ToLiteral().Await()); EXPECT_THAT(received_literal->data(), ElementsAreArray(literal.data())); } TEST(PjRtCApiClientTest, AsyncHostToDeviceTransferManagerTransferLiteral) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); xla::Shape shape = xla::ShapeUtil::MakeShapeWithType({4}); std::vector data = {1, 2, 3, 4}; xla::Literal literal = xla::LiteralUtil::CreateR1(data); std::vector host_shapes = {shape}; - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( std::unique_ptr transfer_manager, client->CreateBuffersForAsyncHostToDevice(absl::MakeSpan(host_shapes), @@ -936,8 +928,8 @@ TEST(PjRtCApiClientTest, AsyncHostToDeviceTransferManagerTransferLiteral) { std::unique_ptr buffer = transfer_manager->RetrieveBuffer(/*buffer_index=*/0); - TF_ASSERT_OK_AND_ASSIGN(std::shared_ptr result_literal, - buffer->ToLiteral().Await()); + ASSERT_OK_AND_ASSIGN(std::shared_ptr result_literal, + buffer->ToLiteral().Await()); EXPECT_TRUE(LiteralTestUtil::Equal(literal, *result_literal)); } @@ -1038,8 +1030,8 @@ ENTRY Identity() -> f32[2, 2] { TEST(PjRtCApiClientTest, AddressableDeviceLogicalIds) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); ASSERT_GT(client->addressable_devices().size(), 1); XlaBuilder builder("Identity"); @@ -1054,8 +1046,8 @@ TEST(PjRtCApiClientTest, AddressableDeviceLogicalIds) { CompileOptions options; options.executable_build_options.set_device_assignment(device_assignment); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, - client->CompileAndLoad(computation, options)); + ASSERT_OK_AND_ASSIGN(std::unique_ptr executable, + client->CompileAndLoad(computation, options)); absl::Span logical_ids = executable->addressable_device_logical_ids(); @@ -1068,20 +1060,20 @@ TEST(PjRtCApiClientTest, AddressableDeviceLogicalIds) { TEST(PjRtCApiClientTest, Bitcast) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); std::vector data{3}; Shape shape = ShapeUtil::MakeShape(S32, {}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( *shape.mutable_layout(), client->GetDefaultLayout(shape.element_type(), shape.dimensions())); Shape new_shape = ShapeUtil::MakeShape(S32, {1}); - TF_ASSERT_OK_AND_ASSIGN(*new_shape.mutable_layout(), - client->GetDefaultLayout(new_shape.element_type(), - new_shape.dimensions())); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN(*new_shape.mutable_layout(), + client->GetDefaultLayout(new_shape.element_type(), + new_shape.dimensions())); + ASSERT_OK_AND_ASSIGN( auto buffer, client->BufferFromHostBuffer( data.data(), shape.element_type(), shape.dimensions(), @@ -1089,7 +1081,7 @@ TEST(PjRtCApiClientTest, Bitcast) { PjRtClient::HostBufferSemantics::kImmutableOnlyDuringCall, nullptr, client->memory_spaces()[0], /*device_layout=*/nullptr)); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( auto bitcast_buffer, buffer->Bitcast(new_shape.element_type(), new_shape.dimensions(), &new_shape.layout())); @@ -1097,29 +1089,29 @@ TEST(PjRtCApiClientTest, Bitcast) { ASSERT_EQ(bitcast_buffer->on_device_shape(), new_shape); auto future = bitcast_buffer->ToLiteral(); - TF_ASSERT_OK_AND_ASSIGN(auto shared_literal, future.Await()); + ASSERT_OK_AND_ASSIGN(auto shared_literal, future.Await()); std::vector expected = {3}; EXPECT_EQ(shared_literal->data(), expected); } TEST(PjRtCApiClientTest, MakeCanonicalShapeForMemorySpace) { SetUpCpuPjRtApi(); - TF_ASSERT_OK_AND_ASSIGN(std::unique_ptr client, - GetCApiClient("cpu")); + ASSERT_OK_AND_ASSIGN(std::unique_ptr client, + GetCApiClient("cpu")); - TF_ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, - client->GetTopologyDescription()); + ASSERT_OK_AND_ASSIGN(const PjRtTopologyDescription* topology, + client->GetTopologyDescription()); ASSERT_NE(topology, nullptr); Shape input_shape = ShapeUtil::MakeShape(F32, {10, 20}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( Shape canonical_shape, topology->MakeCanonicalShapeForMemorySpace( /*memory_space_kind_id=*/0, input_shape, /*layout=*/nullptr)); EXPECT_TRUE(canonical_shape.has_layout()); Layout specific_layout = LayoutUtil::MakeLayout({0, 1}); - TF_ASSERT_OK_AND_ASSIGN( + ASSERT_OK_AND_ASSIGN( Shape canonical_shape_specific, topology->MakeCanonicalShapeForMemorySpace( /*memory_space_kind_id=*/0, input_shape, &specific_layout)); From 3cf2b224f48cf6643897dce6d6595f77135d0095 Mon Sep 17 00:00:00 2001 From: "A. Unique TensorFlower" Date: Tue, 1 Sep 2026 14:41:26 -0700 Subject: [PATCH 40/43] Integrate LLVM at llvm/llvm-project@cbc5a226cbf8 Updates LLVM usage to match [cbc5a226cbf8](https://github.com/llvm/llvm-project/commit/cbc5a226cbf8) PiperOrigin-RevId: 974715562 --- third_party/xla/third_party/llvm/build.patch | 13 - .../xla/third_party/llvm/generated.patch | 1433 +------- .../xla/third_party/llvm/workspace.bzl | 4 +- .../xla/third_party/shardy/temporary.patch | 3027 ++++++++--------- .../xla/third_party/shardy/workspace.bzl | 4 +- .../triton/common/llvm_cl974500093.patch | 198 ++ .../xla/third_party/triton/common/series.bzl | 1 + .../tests/elementwise/convert_s8_s32.hlo | 2 +- .../tiled/tests/elementwise/is_finite.hlo | 2 +- .../tests/elementwise/reduce_precision.hlo | 2 +- .../tiled/tests/transpose/transpose_c64.hlo | 2 +- .../transforms/tests/vectorize_xtile.mlir | 6 +- .../tests/promote_shuffle_to_dpp.mlir | 10 +- .../tests/lower_xla_intrinsic_lib.mlir | 20 +- .../dialects/tests/memref/dim.mlir | 2 +- .../tests/vector/extract_strided_slice.mlir | 4 +- .../tests/vector/insert_strided_slice.mlir | 4 +- .../Dialect/deallocation/buffer_reuse.mlir | 4 +- .../xla/xla/mlir_hlo/tests/alloc_to_arg.mlir | 2 +- .../collapse_parallel_loops_to_1d_pass.mlir | 2 +- .../mlir_hlo/tests/naive_copy_removal.mlir | 14 +- .../xla/xla/mlir_hlo/tests/tile_loops.mlir | 4 +- .../xla/mlir_hlo/tests/vectorize_copy.mlir | 2 +- 23 files changed, 1728 insertions(+), 3034 deletions(-) create mode 100644 third_party/xla/third_party/triton/common/llvm_cl974500093.patch diff --git a/third_party/xla/third_party/llvm/build.patch b/third_party/xla/third_party/llvm/build.patch index e98a6a218fa61e..868226d56a92b7 100644 --- a/third_party/xla/third_party/llvm/build.patch +++ b/third_party/xla/third_party/llvm/build.patch @@ -70,16 +70,3 @@ index a7e652c..5b8ac5e 100644 "//conditions:default": [ "BLAKE3_NO_AVX2", "BLAKE3_NO_AVX512", - -diff --git a/utils/bazel/llvm-project-overlay/llvm/config.bzl b/utils/bazel/llvm-project-overlay/llvm/config.bzl ---- a/utils/bazel/llvm-project-overlay/llvm/config.bzl -+++ b/utils/bazel/llvm-project-overlay/llvm/config.bzl -@@ -110,6 +110,7 @@ - # LLVM features - r'LTDL_SHLIB_EXT=\".dll\"', - r'LLVM_PLUGIN_EXT=\".dll\"', -+ "LLVM_ENABLE_THREADS=1", - ] + fenv_defines - - # TODO: We should switch to platforms-based config settings to make this easier - diff --git a/third_party/xla/third_party/llvm/generated.patch b/third_party/xla/third_party/llvm/generated.patch index 6258925af88cc2..a490f07bce52c9 100644 --- a/third_party/xla/third_party/llvm/generated.patch +++ b/third_party/xla/third_party/llvm/generated.patch @@ -1,65 +1,4 @@ Auto generated patch. Do not edit or delete it, even if empty. -diff -ruN --strip-trailing-cr a/bolt/lib/Core/Relocation.cpp b/bolt/lib/Core/Relocation.cpp ---- a/bolt/lib/Core/Relocation.cpp -+++ b/bolt/lib/Core/Relocation.cpp -@@ -606,14 +606,6 @@ - case ELF::R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC: - case ELF::R_AARCH64_TLSLE_ADD_TPREL_HI12: - case ELF::R_AARCH64_TLSLE_ADD_TPREL_LO12_NC: -- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12: -- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12_NC: -- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12: -- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12_NC: -- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12: -- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC: -- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12: -- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC: - case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0: - case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0_NC: - case ELF::R_AARCH64_TLSDESC_LD64_LO12: -diff -ruN --strip-trailing-cr a/bolt/test/AArch64/tls.c b/bolt/test/AArch64/tls.c ---- a/bolt/test/AArch64/tls.c -+++ b/bolt/test/AArch64/tls.c -@@ -5,8 +5,6 @@ - int b; - } tbssstruct = {}, tdatastruct = {4, 2}; - --__thread int directaccess; -- - extern __thread struct str extstruct; - - extern void processAddr(volatile void *); -@@ -20,9 +18,6 @@ - processAddr(&tbssstruct.b); - processAddr(&tdatastruct.b); - -- // R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC for a direct access -- directaccess++; -- - // The R_AARCH64_TLSIE_ADR_GOTTPREL_PAGE21 and - // R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC relocations - processAddr(&extstruct.b); -@@ -33,8 +28,6 @@ - // RUN: -Wl,--unresolved-symbols=ignore-all \ - // RUN: -fuse-ld=lld \ - // RUN: -nostdlib --// RUN: llvm-objdump -d -r --disassemble-symbols=main %t.exe \ --// RUN: | FileCheck %s --check-prefix=CHECK-DIRECT-ACCESS - // RUN: llvm-bolt %t.exe -o %t.bolt - // RUN: %clang %cflags -fPIC -pie %s -o %t_pie.exe -Wl,-q \ - // RUN: -Wl,--unresolved-symbols=ignore-all \ -@@ -47,11 +40,6 @@ - // RUN: llvm-objdump -d -r --disassemble-symbols=main %t.so | FileCheck %s - // RUN: llvm-bolt %t.so -o %t.bolt.so - --// CHECK-DIRECT-ACCESS: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess --// CHECK-DIRECT-ACCESS-NEXT: add {{.*}} #0x1 --// CHECK-DIRECT-ACCESS-NEXT: str {{.*}} --// CHECK-DIRECT-ACCESS-NEXT: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess -- - // Verify that unoptimized TLS access was generated for shared object. - // CHECK: adrp x0 - // CHECK-NEXT: R_AARCH64_TLSDESC_ADR_PAGE21 tbssstruct diff -ruN --strip-trailing-cr a/lldb/include/lldb/Symbol/Symbol.h b/lldb/include/lldb/Symbol/Symbol.h --- a/lldb/include/lldb/Symbol/Symbol.h +++ b/lldb/include/lldb/Symbol/Symbol.h @@ -90,41 +29,10 @@ diff -ruN --strip-trailing-cr a/lldb/include/lldb/Symbol/Symbol.h b/lldb/include private: union { // Contains the value, or the section offset address when the value is an -diff -ruN --strip-trailing-cr a/lldb/source/API/SBThread.cpp b/lldb/source/API/SBThread.cpp ---- a/lldb/source/API/SBThread.cpp -+++ b/lldb/source/API/SBThread.cpp -@@ -465,7 +465,7 @@ - - // Release the run lock but keep the API lock. - TargetAPIMutex api_mutex = exe_ctx.AllowResume(); -- std::lock_guard guard(api_mutex, std::adopt_lock); -+ std::unique_lock guard(api_mutex, std::adopt_lock); - if (process->GetTarget().GetDebugger().GetAsyncExecution()) - return process->Resume(); - return process->ResumeSynchronous(nullptr); -diff -ruN --strip-trailing-cr a/lldb/unittests/Target/TargetAPIMutexTest.cpp b/lldb/unittests/Target/TargetAPIMutexTest.cpp ---- a/lldb/unittests/Target/TargetAPIMutexTest.cpp -+++ b/lldb/unittests/Target/TargetAPIMutexTest.cpp -@@ -120,6 +120,15 @@ - EXPECT_FALSE(background_lock.try_lock()); - }); - t.join(); -+ -+ // Unlock the original locked mutex. -+ // Calling try_lock() resolves the underlying mutex and re-enters it on this -+ // thread (incrementing the recursive count), so we unlock twice to fully -+ // release both acquisitions. -+ TargetAPIMutex cleanup_lock(target_sp); -+ ASSERT_TRUE(cleanup_lock.try_lock()); -+ cleanup_lock.unlock(); -+ cleanup_lock.unlock(); - } - - TEST_F(TargetAPIMutexTargetTest, LockGuardReleasesOnScopeExit) { diff -ruN --strip-trailing-cr a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp --- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp +++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp -@@ -5937,6 +5937,8 @@ +@@ -5964,6 +5964,8 @@ ISD::matchUnaryPredicate( Y, [&](auto *C) { @@ -133,72 +41,6 @@ diff -ruN --strip-trailing-cr a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/ const APInt &YConst = C->getAsAPIntVal(); return (Opcode == ISD::ABDS) ? YConst.isSignedIntN(Bits) -diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp ---- a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp -+++ b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp -@@ -1254,18 +1254,6 @@ - return true; - } - --/// Check whether \p GAN is the low part of a TLS address computation, i.e. the --/// second operand of an ADDlow. The target flags on their own do not tell the --/// ELF local-exec (:tprel_lo12: and :tprel_lo12_nc:) cases apart from other --/// uses, so callers that depend on local-exec semantics have to check the --/// object format as well. Local dynamic never gets here because it does not --/// build an ADDlow. --static bool isTLSLo12(const GlobalAddressSDNode *GAN) { -- unsigned Flags = GAN->getTargetFlags(); -- return (Flags & (AArch64II::MO_TLS | AArch64II::MO_FRAGMENT)) == -- (AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); --} -- - /// Check if the immediate offset is valid as a scaled immediate. - static bool isValidAsScaledImmediate(int64_t Offset, unsigned Range, - unsigned Size) { -@@ -1361,13 +1349,8 @@ - if (!GAN) - return true; - -- // Folding the low part of an ELF local-exec TLS address into a 128-bit -- // access needs R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC variant, which -- // the GNU bfd linker does not support, so keep materialising the address -- // with an add. - if (GAN->getOffset() % Size == 0 && -- GAN->getGlobal()->getPointerAlignment(DL) >= Size && -- !(Size > 8 && Subtarget->isTargetELF() && isTLSLo12(GAN))) -+ GAN->getGlobal()->getPointerAlignment(DL) >= Size) - return true; - } - -diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp ---- a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp -+++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp -@@ -11542,7 +11542,10 @@ - // add x0, x0, :tprel_lo12:a - SDValue Var = DAG.getTargetGlobalAddress( - GV, DL, PtrVT, 0, AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); -- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, ThreadBase, Var); -+ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, ThreadBase, -+ Var, -+ DAG.getTargetConstant(0, DL, MVT::i32)), -+ 0); - } - - case 24: { -@@ -11558,10 +11561,9 @@ - HiVar, - DAG.getTargetConstant(0, DL, MVT::i32)), - 0); -- // Emit the low part as an ADDlow so that it can be folded into the -- // addressing mode of a following load or store, turning the add into a -- // :tprel_lo12_nc: relocation on the memory access itself. -- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, Addr, LoVar); -+ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, Addr, LoVar, -+ DAG.getTargetConstant(0, DL, MVT::i32)), -+ 0); - } - - case 32: { diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/abd-combine.ll b/llvm/test/CodeGen/AArch64/abd-combine.ll --- a/llvm/test/CodeGen/AArch64/abd-combine.ll +++ b/llvm/test/CodeGen/AArch64/abd-combine.ll @@ -229,1276 +71,3 @@ diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/abd-combine.ll b/llvm/ declare <8 x i16> @llvm.aarch64.neon.sabd.v8i16(<8 x i16>, <8 x i16>) declare <8 x i32> @llvm.abs.v8i32(<8 x i32>, i1) +declare <4 x i32> @llvm.abs.v4i32(<4 x i32>, i1) -diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll ---- a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll -+++ b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll -@@ -27,24 +27,24 @@ - ; RUN: llc -mtriple=arm64-none-linux-gnu -filetype=obj < %s -code-model=large | llvm-objdump -r - | FileCheck --check-prefix=CHECK-24-RELOC %s - - @local_exec_var = thread_local(localexec) global i32 0 --@local_exec_var64 = thread_local(localexec) global i64 0 --@vec_local_exec_var = thread_local(localexec) global <2 x i64> zeroinitializer, align 16 - - define i32 @test_local_exec() { - ; CHECK-LABEL: test_local_exec: - %val = load i32, ptr @local_exec_var - - ; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 --; CHECK-12: ldr w0, [x[[R1]], :tprel_lo12:local_exec_var] -+; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:local_exec_var -+; CHECK-12: ldr w0, [x[[R2]]] - --; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12 -+; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 - - ; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 - ; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var --; CHECK-24: ldr w0, [x[[R2]], :tprel_lo12_nc:local_exec_var] -+; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:local_exec_var -+; CHECK-24: ldr w0, [x[[R3]]] - - ; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 --; CHECK-24-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC -+; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC - - ; CHECK-32: movz x[[R2:[0-9]+]], #:tprel_g1:local_exec_var - ; CHECK-32: mrs x[[R1:[0-9]+]], TPIDR_EL0 -@@ -66,24 +66,6 @@ - ret i32 %val - } - --define void @test_local_exec_store64(i64 %val) { --; CHECK-LABEL: test_local_exec_store64: -- store i64 %val, ptr @local_exec_var64 -- --; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 --; CHECK-12: str x0, [x[[R1]], :tprel_lo12:local_exec_var64] -- --; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST64_TPREL_LO12 -- --; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 --; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var64 --; CHECK-24: str x0, [x[[R2]], :tprel_lo12_nc:local_exec_var64] -- --; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 local_exec_var64 --; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC local_exec_var64 -- ret void --} -- - define ptr @test_local_exec_addr() { - ; CHECK-LABEL: test_local_exec_addr: - ret ptr @local_exec_var -@@ -122,26 +104,3 @@ - ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G1_NC - ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G0_NC - } -- --; A 128-bit access would need R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC --; variant, which not every linker implements, so the low part stays in a --; separate add. --define <2 x i64> @test_local_exec_128bit() { --; CHECK-LABEL: test_local_exec_128bit: -- %val = load <2 x i64>, ptr @vec_local_exec_var -- --; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 --; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:vec_local_exec_var --; CHECK-12: ldr q0, [x[[R2]]] -- --; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 vec_local_exec_var -- --; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 --; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:vec_local_exec_var --; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:vec_local_exec_var --; CHECK-24: ldr q0, [x[[R3]]] -- --; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 vec_local_exec_var --; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC vec_local_exec_var -- ret <2 x i64> %val --} -diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/win-tls.ll b/llvm/test/CodeGen/AArch64/win-tls.ll ---- a/llvm/test/CodeGen/AArch64/win-tls.ll -+++ b/llvm/test/CodeGen/AArch64/win-tls.ll -@@ -3,7 +3,6 @@ - @tlsVar = thread_local global i32 0 - @tlsVar8 = thread_local global i8 0 - @tlsVar64 = thread_local global i64 0 --@tlsVar128 = thread_local global <2 x i64> zeroinitializer - - define i32 @getVar() { - %1 = load i32, ptr @tlsVar -@@ -29,11 +28,6 @@ - ret i64 %1 - } - --define <2 x i64> @getVar128() { -- %1 = load <2 x i64>, ptr @tlsVar128 -- ret <2 x i64> %1 --} -- - ; CHECK-LABEL: getVar - ; CHECK: adrp [[TLS_INDEX_ADDR:x[0-9]+]], _tls_index - ; CHECK: ldr [[TLS_POINTER:x[0-9]+]], [x18, #88] -@@ -68,7 +62,3 @@ - ; CHECK-LABEL: getVar64 - ; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar64 - ; CHECK: ldr x0, [[[TLS]], :secrel_lo12:tlsVar64] -- --; CHECK-LABEL: getVar128 --; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar128 --; CHECK: ldr q0, [[[TLS]], :secrel_lo12:tlsVar128] -diff -ruN --strip-trailing-cr a/utils/bazel/.bazelrc b/utils/bazel/.bazelrc ---- a/utils/bazel/.bazelrc -+++ b/utils/bazel/.bazelrc -@@ -222,6 +222,19 @@ - build:hermetic-toolchain --copt=-Wno-modules-import-nested-redundant --host_copt=-Wno-modules-import-nested-redundant - - ############################################################################### -+# Options for Emscripten WebAssembly builds. -+############################################################################### -+ -+build:wasm --platforms=@emsdk//:platform_wasm -+# Match LLVM's single-threaded config and shut the runtime down when main returns. -+build:wasm --features=-use_pthreads,exit_runtime -+# Make the default CLI artifact use Node's host filesystem; browser builds override this to 0. -+build:wasm --linkopt=-sNODERAWFS=1 -+ -+# TODO: zstd unconditionally enables pthreads on non-Windows targets. -+build:wasm --@llvm-project//third-party:llvm_enable_zstd=false -+ -+############################################################################### - # Options for continuous integration. - ############################################################################### - -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel -@@ -1957,6 +1957,7 @@ - ":parse", - ":sema", - ":serialization", -+ "//compiler-rt:emutls", - "//llvm:AllTargetsAsmParsers", - "//llvm:AllTargetsCodeGens", - "//llvm:Core", -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel -@@ -165,6 +165,15 @@ - ], - ) - -+cc_library( -+ name = "emutls", -+ srcs = [ -+ "lib/builtins/emutls.c", -+ ], -+ hdrs = glob(["lib/builtins/*.h"]), -+ linkstatic = True, -+) -+ - filegroup( - name = "fuzzer_installed_hdrs", - srcs = glob(["include/fuzzer/*.h"]), -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -@@ -1467,15 +1467,12 @@ - name = "__support_libc_assert", - hdrs = ["src/__support/libc_assert.h"], - deps = [ -- ":__support_integer_to_string", - ":__support_macros_attributes", - ":__support_macros_config", - ":__support_macros_hardening", - ":__support_macros_macro_utils", - ":__support_macros_optimization", - ":__support_macros_properties_os", -- ":__support_osutil_exit_hdrs", -- ":__support_osutil_io", - ], - ) - -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl ---- a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl -+++ b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl -@@ -66,7 +66,7 @@ - copts = copts + _FULL_BUILD_COPTS - - # Temporarily disable full_build tests (currently broken) to unblock CI. -- tags = tags + ["manual", "notap"] -+ tags = tags + ["manual", "nobuildkite", "notap"] - cc_test( - name = name, - local_defines = local_defines + _TEST_DEFINES + LIBC_CONFIGURE_OPTIONS, -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel ---- a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel -+++ b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel -@@ -456,6 +456,7 @@ - "-pthread", - "-ldl", - ], -+ "@platforms//os:emscripten": [], - "//conditions:default": [ - "-pthread", - "-ldl", -@@ -1439,7 +1440,10 @@ - "include/llvm/Analysis/Utils/*.h", - ], - ) + ["include/llvm-c/Analysis.h"], -- copts = llvm_copts + ["-ftrapping-math"], -+ copts = llvm_copts + select({ -+ "@platforms//os:emscripten": [], -+ "//conditions:default": ["-ftrapping-math"], -+ }), - features = ["-parse_headers"], - textual_hdrs = glob([ - "include/llvm/Analysis/*.def", -@@ -4973,6 +4977,7 @@ - # ll scripts rely on symbols from dependent - # libraries being resolvable. - linkopts = select({ -+ "@platforms//os:emscripten": [], - "@platforms//os:macos": [], - "@platforms//os:windows": [], - "//conditions:default": [ -@@ -5548,6 +5553,7 @@ - copts = llvm_copts, - # Make symbols from the standard library dynamically resolvable. - linkopts = select({ -+ "@platforms//os:emscripten": [], - "@platforms//os:macos": [], - "@platforms//os:windows": [], - "//conditions:default": [ -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/config.bzl b/utils/bazel/llvm-project-overlay/llvm/config.bzl ---- a/utils/bazel/llvm-project-overlay/llvm/config.bzl -+++ b/utils/bazel/llvm-project-overlay/llvm/config.bzl -@@ -46,7 +46,28 @@ - "HAVE_UNISTD_H=1", - ] - -+emscripten_defines = [ -+ "LLVM_ON_UNIX=1", -+ r'LTDL_SHLIB_EXT=\".so\"', -+ r'LLVM_PLUGIN_EXT=\".so\"', -+ "LLVM_ENABLE_LLVM_EXPORT_ANNOTATIONS=1", -+ "LLVM_ENABLE_PLUGINS=0", -+ "LLVM_ENABLE_THREADS=0", -+ "HAVE_MALLINFO=1", -+ "HAVE_SETENV_R=1", -+ "HAVE_STRERROR_R=1", -+ "HAVE_SYSEXITS_H=1", -+ "HAVE_SYS_IOCTL_H=1", -+ "HAVE_UNISTD_H=1", -+] -+ -+fenv_defines = [ -+ "HAVE_DECL_FE_ALL_EXCEPT=1", -+ "HAVE_DECL_FE_INEXACT=1", -+] -+ - backtrace_defines = select({ -+ "@platforms//os:emscripten": [], - "@platforms//os:windows": [], - "@llvm//platforms/config:musl": [], - "//conditions:default": [ -@@ -60,14 +81,14 @@ - "//conditions:default": [], - }) - --linux_defines = posix_defines + [ -+linux_defines = posix_defines + fenv_defines + [ - "_GNU_SOURCE", - "HAVE_GETAUXVAL=1", - "HAVE_SBRK=1", - "HAVE_STRUCT_STAT_ST_MTIM_TV_NSEC=1", - ] - --macos_defines = posix_defines + [ -+macos_defines = posix_defines + fenv_defines + [ - "HAVE_MACH_MACH_H=1", - "HAVE_MALLOC_MALLOC_H=1", - "HAVE_MALLOC_ZONE_STATISTICS=1", -@@ -89,12 +110,13 @@ - # LLVM features - r'LTDL_SHLIB_EXT=\".dll\"', - r'LLVM_PLUGIN_EXT=\".dll\"', --] -+] + fenv_defines - - # TODO: We should switch to platforms-based config settings to make this easier - # to express. - os_defines = select({ -- "@platforms//os:freebsd": posix_defines, -+ "@platforms//os:emscripten": emscripten_defines, -+ "@platforms//os:freebsd": posix_defines + fenv_defines, - "@platforms//os:macos": macos_defines, - "@platforms//os:windows": win32_defines, - "//conditions:default": linux_defines, -@@ -117,6 +139,7 @@ - Label("//llvm:linux_ppc64le"): native_arch_defines("PowerPC", "powerpc64le-unknown-linux-gnu"), - Label("//llvm:linux_riscv64"): native_arch_defines("RISCV", "riscv64-unknown-linux-gnu"), - Label("//llvm:linux_s390x"): native_arch_defines("SystemZ", "systemz-unknown-linux_gnu"), -+ "@platforms//os:emscripten": native_arch_defines("WebAssembly", "wasm32-unknown-emscripten"), - "@platforms//os:windows": native_arch_defines("X86", "x86_64-pc-win32"), - "//conditions:default": native_arch_defines("X86", "x86_64-unknown-linux-gnu"), - }) + [ -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h ---- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h -+++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h -@@ -56,11 +56,11 @@ - - /* Define to 1 if you have the declaration of `FE_ALL_EXCEPT', and to 0 if you - don't. */ --#define HAVE_DECL_FE_ALL_EXCEPT 1 -+/* HAVE_DECL_FE_ALL_EXCEPT defined in Bazel */ - - /* Define to 1 if you have the declaration of `FE_INEXACT', and to 0 if you - don't. */ --#define HAVE_DECL_FE_INEXACT 1 -+/* HAVE_DECL_FE_INEXACT defined in Bazel */ - - /* Define to 1 if you have the declaration of `strerror_s', and to 0 if you - don't. */ -diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h ---- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h -+++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h -@@ -28,7 +28,7 @@ - /* LLVM_DEFAULT_TARGET_TRIPLE defined in Bazel */ - - /* Define if threads enabled */ --#define LLVM_ENABLE_THREADS 1 -+/* LLVM_ENABLE_THREADS defined in Bazel */ - - /* Has gcc/MSVC atomic intrinsics */ - #define LLVM_HAS_ATOMICS 1 -diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel b/utils/bazel/MODULE.bazel ---- a/utils/bazel/MODULE.bazel -+++ b/utils/bazel/MODULE.bazel -@@ -30,6 +30,7 @@ - bazel_dep(name = "libpfm", version = "4.13.0", repo_name = "pfm") - bazel_dep(name = "vulkan_headers", version = "1.4.349") - -+bazel_dep(name = "emsdk", version = "6.0.2", dev_dependency = True) - bazel_dep(name = "llvm", version = "0.8.5", dev_dependency = True) - - llvm_repos_extension = use_extension(":extensions.bzl", "llvm_repos_extension") -diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel.lock b/utils/bazel/MODULE.bazel.lock ---- a/utils/bazel/MODULE.bazel.lock -+++ b/utils/bazel/MODULE.bazel.lock -@@ -81,6 +81,8 @@ - "https://bcr.bazel.build/modules/buildozer/7.1.2/source.json": "c9028a501d2db85793a6996205c8de120944f50a0d570438fcae0457a5f9d1f8", - "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/MODULE.bazel": "f6561baff0fc0035c9c1a9e2b0820de106cdb01b37bf5c81276860ccc863e5b2", - "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/source.json": "a8611a2b5577929ad7e1f44ded19dab21a188125a74ac6192d21d283609f280f", -+ "https://bcr.bazel.build/modules/emsdk/6.0.2/MODULE.bazel": "4a3c4195e5f2e0056bc18bf9f8af631c4f720c3e8cea45cfb7247eacf02e27fe", -+ "https://bcr.bazel.build/modules/emsdk/6.0.2/source.json": "53111cbcb9f0971aa14da976bafbb937a1bc92a2580c8881890983cd2d289af0", - "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/MODULE.bazel": "f1b7bb2dd53e8f2ef984b39485ec8a44e9076dda5c4b8efd2fb4c6a6e856a31d", - "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/source.json": "ebe931bfe362e4b41e59ee00a528db6074157ff2ced92eb9e970acab2e1089c9", - "https://bcr.bazel.build/modules/gazelle/0.32.0/MODULE.bazel": "b499f58a5d0d3537f3cf5b76d8ada18242f64ec474d8391247438bf04f58c7b8", -@@ -184,6 +186,7 @@ - "https://bcr.bazel.build/modules/rules_cc/0.2.13/MODULE.bazel": "eecdd666eda6be16a8d9dc15e44b5c75133405e820f620a234acc4b1fdc5aa37", - "https://bcr.bazel.build/modules/rules_cc/0.2.14/MODULE.bazel": "353c99ed148887ee89c54a17d4100ae7e7e436593d104b668476019023b58df8", - "https://bcr.bazel.build/modules/rules_cc/0.2.15/MODULE.bazel": "6a0a4a75a57aa6dc888300d848053a58c6b12a29f89d4304e1c41448514ec6e8", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.16/MODULE.bazel": "9242fa89f950c6ef7702801ab53922e99c69b02310c39fb6e62b2bd30df2a1d4", - "https://bcr.bazel.build/modules/rules_cc/0.2.17/MODULE.bazel": "1849602c86cb60da8613d2de887f9566a6d354a6df6d7009f9d04a14402f9a84", - "https://bcr.bazel.build/modules/rules_cc/0.2.18/MODULE.bazel": "4460ec36adc8f722a6a2a4ac9374cb91f2acebadaa93fc37966129afb3dece87", - "https://bcr.bazel.build/modules/rules_cc/0.2.19/MODULE.bazel": "d5e0f05b63273281a16654eb6b1a8742a75ec153ac8b4f0419949d6e401e46f0", -@@ -242,6 +245,8 @@ - "https://bcr.bazel.build/modules/rules_license/0.0.7/MODULE.bazel": "088fbeb0b6a419005b89cf93fe62d9517c0a2b8bb56af3244af65ecfe37e7d5d", - "https://bcr.bazel.build/modules/rules_license/1.0.0/MODULE.bazel": "a7fda60eefdf3d8c827262ba499957e4df06f659330bbe6cdbdb975b768bb65c", - "https://bcr.bazel.build/modules/rules_license/1.0.0/source.json": "a52c89e54cc311196e478f8382df91c15f7a2bfdf4c6cd0e2675cc2ff0b56efb", -+ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/MODULE.bazel": "c22a48b2a0dbf05a9dc5f83837bbc24c226c1f6e618de3c3a610044c9f336056", -+ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/source.json": "a3f966f4415a8a6545e560ee5449eac95cc633f96429d08e87c87775c72f5e09", - "https://bcr.bazel.build/modules/rules_pkg/0.7.0/MODULE.bazel": "df99f03fc7934a4737122518bb87e667e62d780b610910f0447665a7e2be62dc", - "https://bcr.bazel.build/modules/rules_pkg/1.0.1/MODULE.bazel": "5b1df97dbc29623bccdf2b0dcd0f5cb08e2f2c9050aab1092fd39a41e82686ff", - "https://bcr.bazel.build/modules/rules_pkg/1.0.1/source.json": "bd82e5d7b9ce2d31e380dd9f50c111d678c3bdaca190cb76b0e1c71b05e1ba8a", -@@ -268,7 +273,8 @@ - "https://bcr.bazel.build/modules/rules_python/1.6.0/MODULE.bazel": "7e04ad8f8d5bea40451cf80b1bd8262552aa73f841415d20db96b7241bd027d8", - "https://bcr.bazel.build/modules/rules_python/1.7.0/MODULE.bazel": "d01f995ecd137abf30238ad9ce97f8fc3ac57289c8b24bd0bf53324d937a14f8", - "https://bcr.bazel.build/modules/rules_python/1.8.0/MODULE.bazel": "c151c025dbcc93d8f62ab68ecc313c9176a868a0e6386981bf2a12aec77cbe7b", -- "https://bcr.bazel.build/modules/rules_python/1.8.0/source.json": "356397eed5b46971d8c585c92098d70495078a80bf18bebcb4209f44b495f3e6", -+ "https://bcr.bazel.build/modules/rules_python/1.8.4/MODULE.bazel": "33e3971e66161a3e955f7a0d411a8d1f291c4ce4c561851512466f3c77ff8ece", -+ "https://bcr.bazel.build/modules/rules_python/1.8.4/source.json": "9fbc0e57bae52cddcc3831d668bce87a47e0c655104a85098d4459dd9a3b0a10", - "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/MODULE.bazel": "d44fec647d0aeb67b9f3b980cf68ba634976f3ae7ccd6c07d790b59b87a4f251", - "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/source.json": "37c10335f2361c337c5c1f34ed36d2da70534c23088062b33a8bdaab68aa9dea", - "https://bcr.bazel.build/modules/rules_rust/0.69.0/MODULE.bazel": "4326fec48f2fef0d514de46346f7f77e200c82936dd08b91c9ef039fbdad5c10", -@@ -348,9 +354,142 @@ - ] - } - }, -+ "@@emsdk+//:emscripten_cache.bzl%emscripten_cache": { -+ "general": { -+ "bzlTransitiveDigest": "GMscy7c4sDbvbf9dMSrtUvyJJBuxYpJBT/Lg+2ob6dk=", -+ "usagesDigest": "Id/C4z1d3MUlKZgmBOQiLi2E7NQZTGsd4WV2wSMzmo4=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+ "generatedRepoSpecs": { -+ "emscripten_cache": { -+ "repoRuleId": "@@emsdk+//:emscripten_cache.bzl%_emscripten_cache_repository", -+ "attributes": { -+ "configuration": [], -+ "targets": [], -+ "prebuilt_cache_url": "", -+ "prebuilt_cache_sha256": "", -+ "prebuilt_cache_strip_prefix": "" -+ } -+ } -+ }, -+ "recordedRepoMappingEntries": [] -+ } -+ }, -+ "@@emsdk+//:emscripten_deps.bzl%emscripten_deps": { -+ "general": { -+ "bzlTransitiveDigest": "ZT33Pf8H8gJ/X4gT3oXVPzuAO0H2y1HtHMPOfILGV0M=", -+ "usagesDigest": "4SlUap0Npa9PDUrLoi0uZ7CRDp9h/YbWi0UuidttuWc=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+ "generatedRepoSpecs": { -+ "emscripten_bin_linux": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -+ "sha256": "d574428df9ecf00790e28636bdc47027432737c31621b18cdb418123afda4ac1", -+ "strip_prefix": "install", -+ "type": "tar.xz", -+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" -+ } -+ }, -+ "emscripten_bin_linux_arm64": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -+ "sha256": "d74803ef563511b9cc1e5cde5016f06d161ffd2b6223135a8aeeef44194594e7", -+ "strip_prefix": "install", -+ "type": "tar.xz", -+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" -+ } -+ }, -+ "emscripten_bin_mac": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -+ "sha256": "356f36ba04a54edb029c658dd1b547c5c8a8f3c166b09654c1efcb9cf7bf8a57", -+ "strip_prefix": "install", -+ "type": "tar.xz", -+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" -+ } -+ }, -+ "emscripten_bin_mac_arm64": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -+ "sha256": "ded3bb783e7aa3dda576955dd0aa3a71dd21789e42befb63ee14f7d9f9b6aa32", -+ "strip_prefix": "install", -+ "type": "tar.xz", -+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" -+ } -+ }, -+ "emscripten_bin_win": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/clang++.exe\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/llvm-ar.exe\",\n \"bin/llvm-dwarfdump.exe\",\n \"bin/llvm-nm.exe\",\n \"bin/llvm-objcopy.exe\",\n \"bin/wasm-ctor-eval.exe\",\n \"bin/wasm-emscripten-finalize.exe\",\n \"bin/wasm-ld.exe\",\n \"bin/wasm-metadce.exe\",\n \"bin/wasm-opt.exe\",\n \"bin/wasm-split.exe\",\n \"bin/wasm2js.exe\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar.exe\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -+ "sha256": "e5f9250a9cf4ff6ed16d57d6b5e177c844067381d31cd0c1a607c1ee1d2ba088", -+ "strip_prefix": "install", -+ "type": "zip", -+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/win/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.zip" -+ } -+ } -+ }, -+ "recordedRepoMappingEntries": [ -+ [ -+ "emsdk+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "emsdk+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_cc+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "rules_cc+", -+ "cc_compatibility_proxy", -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -+ ], -+ [ -+ "rules_cc+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -+ "rules_cc", -+ "rules_cc+" -+ ] -+ ] -+ } -+ }, -+ "@@protobuf+//python/dist:system_python.bzl%system_python_extension": { -+ "general": { -+ "bzlTransitiveDigest": "qh0n9IrXU/xS94wxKQrG1J63zrLkA1Wy2Y3BQxptPcI=", -+ "usagesDigest": "tCi55FyqtOJ2jXh9vcjrHCl4ov3kpWiwKl103nA9BOI=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+ "generatedRepoSpecs": { -+ "system_python": { -+ "repoRuleId": "@@protobuf+//python/dist:system_python.bzl%system_python", -+ "attributes": { -+ "minimum_python_version": "3.9" -+ } -+ } -+ }, -+ "recordedRepoMappingEntries": [] -+ } -+ }, - "@@pybind11_bazel+//:internal_configure.bzl%internal_configure_extension": { - "general": { -- "bzlTransitiveDigest": "NFQjcZF+fAvf5fDH+pqsx4JrfzP9PuHBz6S6ZutIbnw=", -+ "bzlTransitiveDigest": "7zBsfo5dyMqKT23rXrvWqJMx0AugwL6NyirkmvzKcqU=", - "usagesDigest": "D1r3lfzMuUBFxgG8V6o0bQTLMk3GkaGOaPzw53wrwyw=", - "recordedFileInputs": { - "@@pybind11_bazel+//MODULE.bazel": "e6f4c20442eaa7c90d7190d8dc539d0ab422f95c65a57cc59562170c58ae3d34" -@@ -380,8 +519,8 @@ - }, - "@@rules_android+//bzlmod_extensions:apksig.bzl%apksig_extension": { - "general": { -- "bzlTransitiveDigest": "By9qVNN7G4oL1vYOJXye7Dp/CbR2ar9oxAW8WXAVcVw=", -- "usagesDigest": "xq6OVkELeJvOgYo3oY/sUBsGFbcqdV+9BYiNgSPV/po=", -+ "bzlTransitiveDigest": "15xx/lo4VYL9KdLW0Cc94ebALMF07+XZH8dZcVU8/LI=", -+ "usagesDigest": "S8lLnnZxdeYUYq3kIGhVMk0wQ9Fd6elmCskvn+SL6iw=", - "recordedFileInputs": {}, - "recordedDirentsInputs": {}, - "envVariables": {}, -@@ -389,7 +528,10 @@ - "apksig": { - "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", - "attributes": { -- "url": "https://android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz", -+ "urls": [ -+ "https://mirror.bazel.build/android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz" -+ ], -+ "sha256": "12e44fdbd219c5e1cc62099c2a01d775957603d2d4f693f8285f9d95d9a04e77", - "build_file": "@@rules_android+//bzlmod_extensions:apksig.BUILD" - } - } -@@ -405,8 +547,8 @@ - }, - "@@rules_android+//bzlmod_extensions:com_android_dex.bzl%com_android_dex_extension": { - "general": { -- "bzlTransitiveDigest": "rvWbJQc8jInfIAaXIMhSOqUlwM9HVeLey6q0ISvg08Y=", -- "usagesDigest": "toF8IFMu98H/VU2p1sfVC5fVXVYJunpbbmtM6tOsQXY=", -+ "bzlTransitiveDigest": "K0jbWcRwfM8njdIXNRjRvdApKmBfKeFLScDH+5LSSE0=", -+ "usagesDigest": "0hluQmaWiWak6sVMP5L4wXhNyIwv9fw0y5JJ8lnPb1c=", - "recordedFileInputs": {}, - "recordedDirentsInputs": {}, - "envVariables": {}, -@@ -414,8 +556,11 @@ - "com_android_dex": { - "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", - "attributes": { -- "url": "https://android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz", -- "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD" -+ "urls": [ -+ "https://mirror.bazel.build/android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz" -+ ], -+ "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD", -+ "sha256": "86b4848c038bf687fadc812239cb01fb8d1d15cef3125b480a0448360992b95d" - } - } - }, -@@ -444,10 +589,144 @@ - "recordedRepoMappingEntries": [] - } - }, -+ "@@rules_nodejs+//nodejs:extensions.bzl%node": { -+ "general": { -+ "bzlTransitiveDigest": "4pUxCNc22K4I+6+4Nxu52Hur12tFRfa1JMsN5mdDv60=", -+ "usagesDigest": "dqOjZvNvw6/DVBPAiKrXJNA0Tx4GT4Vj/VdUyGMpDL8=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+ "generatedRepoSpecs": { -+ "nodejs_linux_amd64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "linux_amd64" -+ } -+ }, -+ "nodejs_linux_arm64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "linux_arm64" -+ } -+ }, -+ "nodejs_linux_s390x": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "linux_s390x" -+ } -+ }, -+ "nodejs_linux_ppc64le": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "linux_ppc64le" -+ } -+ }, -+ "nodejs_darwin_amd64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "darwin_amd64" -+ } -+ }, -+ "nodejs_darwin_arm64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "darwin_arm64" -+ } -+ }, -+ "nodejs_windows_amd64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "windows_amd64" -+ } -+ }, -+ "nodejs_windows_arm64": { -+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -+ "attributes": { -+ "node_download_auth": {}, -+ "node_repositories": {}, -+ "node_urls": [ -+ "https://nodejs.org/dist/v{version}/{filename}" -+ ], -+ "node_version": "20.18.0", -+ "include_headers": false, -+ "platform": "windows_arm64" -+ } -+ }, -+ "nodejs": { -+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", -+ "attributes": { -+ "user_node_repository_name": "nodejs" -+ } -+ }, -+ "nodejs_host": { -+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", -+ "attributes": { -+ "user_node_repository_name": "nodejs" -+ } -+ }, -+ "nodejs_toolchains": { -+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_toolchains_repo.bzl%nodejs_toolchains_repo", -+ "attributes": { -+ "user_node_repository_name": "nodejs" -+ } -+ } -+ }, -+ "recordedRepoMappingEntries": [] -+ } -+ }, - "@@rules_python+//python/extensions:config.bzl%config": { - "general": { - "bzlTransitiveDigest": "EcMcbtKZvYmd5Mi1Fpg4EeBBztLHEE5tjO5tLDBYDuU=", -- "usagesDigest": "EocbSr4I3/Shk4QaFool8b8navUiUFqgzF9bOaaYfFk=", -+ "usagesDigest": "p2al+dDKI5UlCyNvheMVynbWSGbdiji/jMz53fMNfJA=", - "recordedFileInputs": {}, - "recordedDirentsInputs": {}, - "envVariables": {}, -@@ -682,7 +961,7 @@ - "@@rules_python+//python/uv:uv.bzl%uv": { - "general": { - "bzlTransitiveDigest": "ijW9KS7qsIY+yBVvJ+Nr1mzwQox09j13DnE3iIwaeTM=", -- "usagesDigest": "yXvWfXAzpBeW71mWgwU3AqAzXX/dFACnx12eYvBsJ8w=", -+ "usagesDigest": "/HRt5Hw/vpDr9CDrKEPjeDIjxo4307VLxMu8BNAEDWA=", - "recordedFileInputs": {}, - "recordedDirentsInputs": {}, - "envVariables": {}, -@@ -719,6 +998,533 @@ - ] - ] - } -+ }, -+ "@@rules_rust+//crate_universe:extension.bzl%crate": { -+ "general": { -+ "bzlTransitiveDigest": "VVbU93QvGxFMzb9BcpYTYyyYDpj10Ya6Zm5RH1JEUhw=", -+ "usagesDigest": "EuFUqVKVHF263jHTWOHXs4tFACdRNVOhwpoytdk19bs=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": { -+ "CARGO_BAZEL_DEBUG": null, -+ "CARGO_BAZEL_GENERATOR_SHA256": null, -+ "CARGO_BAZEL_GENERATOR_URL": null, -+ "CARGO_BAZEL_ISOLATED": null, -+ "CARGO_BAZEL_REPIN": null, -+ "CARGO_BAZEL_REPIN_ONLY": null, -+ "CARGO_BAZEL_TIMEOUT": null, -+ "REPIN": null -+ }, -+ "generatedRepoSpecs": { -+ "crates": { -+ "repoRuleId": "@@rules_rust+//crate_universe:extensions.bzl%_generate_repo", -+ "attributes": { -+ "contents": { -+ "BUILD.bazel": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\npackage(default_visibility = [\"//visibility:public\"])\n\nexports_files(\n [\n \"cargo-bazel.json\",\n \"crates.bzl\",\n \"defs.bzl\",\n ] + glob(\n allow_empty = True,\n include = [\"*.bazel\"],\n ),\n)\n\nfilegroup(\n name = \"srcs\",\n srcs = glob(\n allow_empty = True,\n include = [\n \"*.bazel\",\n \"*.bzl\",\n ],\n ),\n)\n\n# Workspace Member Dependencies\nalias(\n name = \"googletest-0.14.3\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"googletest\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme-0.3.37\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste-1.0.15\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote-1.0.47\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn-3.0.3\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n", -+ "alias_rules.bzl": "\"\"\"Alias that transitions its target to `compilation_mode=opt`. Use `transition_alias=\"opt\"` to enable.\"\"\"\n\nload(\"@rules_cc//cc:defs.bzl\", \"CcInfo\")\nload(\"@rules_rust//rust:rust_common.bzl\", \"COMMON_PROVIDERS\")\n\ndef _transition_alias_impl(ctx):\n # `ctx.attr.actual` is a list of 1 item due to the transition\n providers = [ctx.attr.actual[0][provider] for provider in COMMON_PROVIDERS]\n if CcInfo in ctx.attr.actual[0]:\n providers.append(ctx.attr.actual[0][CcInfo])\n return providers\n\ndef _change_compilation_mode(compilation_mode):\n def _change_compilation_mode_impl(_settings, _attr):\n return {\n \"//command_line_option:compilation_mode\": compilation_mode,\n }\n\n return transition(\n implementation = _change_compilation_mode_impl,\n inputs = [],\n outputs = [\n \"//command_line_option:compilation_mode\",\n ],\n )\n\ndef _transition_alias_rule(compilation_mode):\n return rule(\n implementation = _transition_alias_impl,\n provides = COMMON_PROVIDERS,\n attrs = {\n \"actual\": attr.label(\n mandatory = True,\n doc = \"`rust_library()` target to transition to `compilation_mode=opt`.\",\n providers = COMMON_PROVIDERS,\n cfg = _change_compilation_mode(compilation_mode),\n ),\n \"_allowlist_function_transition\": attr.label(\n default = \"@bazel_tools//tools/allowlists/function_transition_allowlist\",\n ),\n },\n doc = \"Transitions a Rust library crate to the `compilation_mode=opt`.\",\n )\n\ntransition_alias_dbg = _transition_alias_rule(\"dbg\")\ntransition_alias_fastbuild = _transition_alias_rule(\"fastbuild\")\ntransition_alias_opt = _transition_alias_rule(\"opt\")\n", -+ "defs.bzl": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\"\"\"\n# `crates_repository` API\n\n- [aliases](#aliases)\n- [crate_deps](#crate_deps)\n- [all_crate_deps](#all_crate_deps)\n- [crate_repositories](#crate_repositories)\n\n\"\"\"\n\nload(\"@bazel_tools//tools/build_defs/repo:git.bzl\", \"new_git_repository\")\nload(\"@bazel_tools//tools/build_defs/repo:http.bzl\", \"http_archive\")\nload(\"@bazel_tools//tools/build_defs/repo:utils.bzl\", \"maybe\")\nload(\"@bazel_skylib//lib:selects.bzl\", \"selects\")\nload(\"@rules_rust//crate_universe/private:local_crate_mirror.bzl\", \"local_crate_mirror\")\n\n###############################################################################\n# MACROS API\n###############################################################################\n\n# An identifier that represent common dependencies (unconditional).\n_COMMON_CONDITION = \"\"\n\ndef _flatten_dependency_maps(all_dependency_maps):\n \"\"\"Flatten a list of dependency maps into one dictionary.\n\n Dependency maps have the following structure:\n\n ```python\n DEPENDENCIES_MAP = {\n # The first key in the map is a Bazel package\n # name of the workspace this file is defined in.\n \"workspace_member_package\": {\n\n # Not all dependencies are supported for all platforms.\n # the condition key is the condition required to be true\n # on the host platform.\n \"condition\": {\n\n # An alias to a crate target. # The label of the crate target the\n # Aliases are only crate names. # package name refers to.\n \"package_name\": \"@full//:label\",\n }\n }\n }\n ```\n\n Args:\n all_dependency_maps (list): A list of dicts as described above\n\n Returns:\n dict: A dictionary as described above\n \"\"\"\n dependencies = {}\n\n for workspace_deps_map in all_dependency_maps:\n for pkg_name, conditional_deps_map in workspace_deps_map.items():\n if pkg_name not in dependencies:\n non_frozen_map = dict()\n for key, values in conditional_deps_map.items():\n non_frozen_map.update({key: dict(values.items())})\n dependencies.setdefault(pkg_name, non_frozen_map)\n continue\n\n for condition, deps_map in conditional_deps_map.items():\n # If the condition has not been recorded, do so and continue\n if condition not in dependencies[pkg_name]:\n dependencies[pkg_name].setdefault(condition, dict(deps_map.items()))\n continue\n\n # Alert on any miss-matched dependencies\n inconsistent_entries = []\n for crate_name, crate_label in deps_map.items():\n existing = dependencies[pkg_name][condition].get(crate_name)\n if existing and existing != crate_label:\n inconsistent_entries.append((crate_name, existing, crate_label))\n dependencies[pkg_name][condition].update({crate_name: crate_label})\n\n return dependencies\n\ndef crate_deps(deps, package_name = None):\n \"\"\"Finds the fully qualified label of the requested crates for the package where this macro is called.\n\n Args:\n deps (list): The desired list of crate targets.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()`.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if not deps:\n return []\n\n if package_name == None:\n package_name = native.package_name()\n\n # Join both sets of dependencies\n dependencies = _flatten_dependency_maps([\n _NORMAL_DEPENDENCIES,\n _NORMAL_DEV_DEPENDENCIES,\n _PROC_MACRO_DEPENDENCIES,\n _PROC_MACRO_DEV_DEPENDENCIES,\n _BUILD_DEPENDENCIES,\n _BUILD_PROC_MACRO_DEPENDENCIES,\n ]).pop(package_name, {})\n\n # Combine all conditional packages so we can easily index over a flat list\n # TODO: Perhaps this should actually return select statements and maintain\n # the conditionals of the dependencies\n flat_deps = {}\n for deps_set in dependencies.values():\n for crate_name, crate_label in deps_set.items():\n flat_deps.update({crate_name: crate_label})\n\n missing_crates = []\n crate_targets = []\n for crate_target in deps:\n if crate_target not in flat_deps:\n missing_crates.append(crate_target)\n else:\n crate_targets.append(flat_deps[crate_target])\n\n if missing_crates:\n fail(\"Could not find crates `{}` among dependencies of `{}`. Available dependencies were `{}`\".format(\n missing_crates,\n package_name,\n dependencies,\n ))\n\n return crate_targets\n\ndef all_crate_deps(\n normal = False, \n normal_dev = False, \n proc_macro = False, \n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Finds the fully qualified label of all requested direct crate dependencies \\\n for the package where this macro is called.\n\n If no parameters are set, all normal dependencies are returned. Setting any one flag will\n otherwise impact the contents of the returned list.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list.\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_dependency_maps = []\n if normal:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n if normal_dev:\n all_dependency_maps.append(_NORMAL_DEV_DEPENDENCIES)\n if proc_macro:\n all_dependency_maps.append(_PROC_MACRO_DEPENDENCIES)\n if proc_macro_dev:\n all_dependency_maps.append(_PROC_MACRO_DEV_DEPENDENCIES)\n if build:\n all_dependency_maps.append(_BUILD_DEPENDENCIES)\n if build_proc_macro:\n all_dependency_maps.append(_BUILD_PROC_MACRO_DEPENDENCIES)\n\n # Default to always using normal dependencies\n if not all_dependency_maps:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n\n dependencies = _flatten_dependency_maps(all_dependency_maps).pop(package_name, None)\n\n if not dependencies:\n if dependencies == None:\n fail(\"Tried to get all_crate_deps for package \" + package_name + \" but that package had no Cargo.toml file\")\n else:\n return []\n\n crate_deps = list(dependencies.pop(_COMMON_CONDITION, {}).values())\n for condition, deps in dependencies.items():\n crate_deps += selects.with_or({\n tuple(_CONDITIONS[condition]): deps.values(),\n \"//conditions:default\": [],\n })\n\n return crate_deps\n\ndef aliases(\n normal = False,\n normal_dev = False,\n proc_macro = False,\n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Produces a map of Crate alias names to their original label\n\n If no dependency kinds are specified, `normal` and `proc_macro` are used by default.\n Setting any one flag will otherwise determine the contents of the returned dict.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list..\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n dict: The aliases of all associated packages\n \"\"\"\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_aliases_maps = []\n if normal:\n all_aliases_maps.append(_NORMAL_ALIASES)\n if normal_dev:\n all_aliases_maps.append(_NORMAL_DEV_ALIASES)\n if proc_macro:\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n if proc_macro_dev:\n all_aliases_maps.append(_PROC_MACRO_DEV_ALIASES)\n if build:\n all_aliases_maps.append(_BUILD_ALIASES)\n if build_proc_macro:\n all_aliases_maps.append(_BUILD_PROC_MACRO_ALIASES)\n\n # Default to always using normal aliases\n if not all_aliases_maps:\n all_aliases_maps.append(_NORMAL_ALIASES)\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n\n aliases = _flatten_dependency_maps(all_aliases_maps).pop(package_name, None)\n\n if not aliases:\n return dict()\n\n common_items = aliases.pop(_COMMON_CONDITION, {}).items()\n\n # If there are only common items in the dictionary, immediately return them\n if not len(aliases.keys()) == 1:\n return dict(common_items)\n\n # Build a single select statement where each conditional has accounted for the\n # common set of aliases.\n crate_aliases = {\"//conditions:default\": dict(common_items)}\n for condition, deps in aliases.items():\n condition_triples = _CONDITIONS[condition]\n for triple in condition_triples:\n if triple in crate_aliases:\n crate_aliases[triple].update(deps)\n else:\n crate_aliases.update({triple: dict(deps.items() + common_items)})\n\n return select(crate_aliases)\n\n###############################################################################\n# WORKSPACE MEMBER DEPS AND ALIASES\n###############################################################################\n\n_NORMAL_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"googletest\": Label(\"@crates//:googletest-0.14.3\"),\n \"linkme\": Label(\"@crates//:linkme-0.3.37\"),\n \"quote\": Label(\"@crates//:quote-1.0.47\"),\n \"syn\": Label(\"@crates//:syn-3.0.3\"),\n },\n },\n}\n\n\n_NORMAL_ALIASES = {\n \"\": {\n _COMMON_CONDITION: {\n },\n },\n}\n\n\n_NORMAL_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_NORMAL_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"paste\": Label(\"@crates//:paste-1.0.15\"),\n },\n },\n}\n\n\n_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_CONDITIONS = {\n \"aarch64-apple-darwin\": [\"@rules_rust//rust/platform:aarch64-apple-darwin\"],\n \"aarch64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\"],\n \"wasm32-unknown-unknown\": [\"@rules_rust//rust/platform:wasm32-unknown-unknown\"],\n \"wasm32-wasip1\": [\"@rules_rust//rust/platform:wasm32-wasip1\"],\n \"x86_64-pc-windows-msvc\": [\"@rules_rust//rust/platform:x86_64-pc-windows-msvc\"],\n \"x86_64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\"],\n \"x86_64-unknown-nixos-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\"],\n}\n\n###############################################################################\n\ndef crate_repositories():\n \"\"\"A macro for defining repositories for all generated crates.\n\n Returns:\n A list of repos visible to the module through the module extension.\n \"\"\"\n maybe(\n http_archive,\n name = \"crates__aho-corasick-1.1.5\",\n sha256 = \"c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/aho-corasick/1.1.5/download\"],\n strip_prefix = \"aho-corasick-1.1.5\",\n build_file = Label(\"@crates//crates:BUILD.aho-corasick-1.1.5.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__autocfg-1.5.1\",\n sha256 = \"f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/autocfg/1.5.1/download\"],\n strip_prefix = \"autocfg-1.5.1\",\n build_file = Label(\"@crates//crates:BUILD.autocfg-1.5.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest-0.14.3\",\n sha256 = \"f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest/0.14.3/download\"],\n strip_prefix = \"googletest-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest_macro-0.14.3\",\n sha256 = \"2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest_macro/0.14.3/download\"],\n strip_prefix = \"googletest_macro-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest_macro-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-0.3.37\",\n sha256 = \"3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme/0.3.37/download\"],\n strip_prefix = \"linkme-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-impl-0.3.37\",\n sha256 = \"77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme-impl/0.3.37/download\"],\n strip_prefix = \"linkme-impl-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-impl-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__memchr-2.8.3\",\n sha256 = \"cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/memchr/2.8.3/download\"],\n strip_prefix = \"memchr-2.8.3\",\n build_file = Label(\"@crates//crates:BUILD.memchr-2.8.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__num-traits-0.2.19\",\n sha256 = \"071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/num-traits/0.2.19/download\"],\n strip_prefix = \"num-traits-0.2.19\",\n build_file = Label(\"@crates//crates:BUILD.num-traits-0.2.19.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__paste-1.0.15\",\n sha256 = \"57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/paste/1.0.15/download\"],\n strip_prefix = \"paste-1.0.15\",\n build_file = Label(\"@crates//crates:BUILD.paste-1.0.15.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__proc-macro2-1.0.107\",\n sha256 = \"985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/proc-macro2/1.0.107/download\"],\n strip_prefix = \"proc-macro2-1.0.107\",\n build_file = Label(\"@crates//crates:BUILD.proc-macro2-1.0.107.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__quote-1.0.47\",\n sha256 = \"1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/quote/1.0.47/download\"],\n strip_prefix = \"quote-1.0.47\",\n build_file = Label(\"@crates//crates:BUILD.quote-1.0.47.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-1.13.1\",\n sha256 = \"f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex/1.13.1/download\"],\n strip_prefix = \"regex-1.13.1\",\n build_file = Label(\"@crates//crates:BUILD.regex-1.13.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-automata-0.4.18\",\n sha256 = \"ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-automata/0.4.18/download\"],\n strip_prefix = \"regex-automata-0.4.18\",\n build_file = Label(\"@crates//crates:BUILD.regex-automata-0.4.18.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-syntax-0.8.11\",\n sha256 = \"d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-syntax/0.8.11/download\"],\n strip_prefix = \"regex-syntax-0.8.11\",\n build_file = Label(\"@crates//crates:BUILD.regex-syntax-0.8.11.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__rustversion-1.0.23\",\n sha256 = \"cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/rustversion/1.0.23/download\"],\n strip_prefix = \"rustversion-1.0.23\",\n build_file = Label(\"@crates//crates:BUILD.rustversion-1.0.23.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-2.0.119\",\n sha256 = \"872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/2.0.119/download\"],\n strip_prefix = \"syn-2.0.119\",\n build_file = Label(\"@crates//crates:BUILD.syn-2.0.119.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-3.0.3\",\n sha256 = \"53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/3.0.3/download\"],\n strip_prefix = \"syn-3.0.3\",\n build_file = Label(\"@crates//crates:BUILD.syn-3.0.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__unicode-ident-1.0.24\",\n sha256 = \"e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/unicode-ident/1.0.24/download\"],\n strip_prefix = \"unicode-ident-1.0.24\",\n build_file = Label(\"@crates//crates:BUILD.unicode-ident-1.0.24.bazel\"),\n )\n\n return [\n struct(repo=\"crates__googletest-0.14.3\", is_dev_dep = False),\n struct(repo=\"crates__linkme-0.3.37\", is_dev_dep = False),\n struct(repo=\"crates__paste-1.0.15\", is_dev_dep = False),\n struct(repo=\"crates__quote-1.0.47\", is_dev_dep = False),\n struct(repo=\"crates__syn-3.0.3\", is_dev_dep = False),\n ]\n" -+ } -+ } -+ }, -+ "crates__aho-corasick-1.1.5": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/aho-corasick/1.1.5/download" -+ ], -+ "strip_prefix": "aho-corasick-1.1.5", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"aho_corasick\",\n deps = [\n \"@crates__memchr-2.8.3//:memchr\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"perf-literal\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=aho-corasick\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.1.5\",\n)\n" -+ } -+ }, -+ "crates__autocfg-1.5.1": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/autocfg/1.5.1/download" -+ ], -+ "strip_prefix": "autocfg-1.5.1", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"autocfg\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2015\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=autocfg\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.5.1\",\n)\n" -+ } -+ }, -+ "crates__googletest-0.14.3": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/googletest/0.14.3/download" -+ ], -+ "strip_prefix": "googletest-0.14.3", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"googletest\",\n deps = [\n \"@crates__num-traits-0.2.19//:num_traits\",\n \"@crates__regex-1.13.1//:regex\",\n ],\n proc_macro_deps = [\n \"@crates__googletest_macro-0.14.3//:googletest_macro\",\n \"@crates__rustversion-1.0.23//:rustversion\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" -+ } -+ }, -+ "crates__googletest_macro-0.14.3": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/googletest_macro/0.14.3/download" -+ ], -+ "strip_prefix": "googletest_macro-0.14.3", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"googletest_macro\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-2.0.119//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest_macro\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" -+ } -+ }, -+ "crates__linkme-0.3.37": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/linkme/0.3.37/download" -+ ], -+ "strip_prefix": "linkme-0.3.37", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"linkme\",\n deps = [\n \"@crates__linkme-0.3.37//:build_script_build\",\n ],\n proc_macro_deps = [\n \"@crates__linkme-impl-0.3.37//:linkme_impl\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__linkme-impl-0.3.37": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/linkme-impl/0.3.37/download" -+ ], -+ "strip_prefix": "linkme-impl-0.3.37", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"linkme_impl\",\n deps = [\n \"@crates__linkme-impl-0.3.37//:build_script_build\",\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-3.0.3//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme-impl\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__memchr-2.8.3": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/memchr/2.8.3/download" -+ ], -+ "strip_prefix": "memchr-2.8.3", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"memchr\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=memchr\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.8.3\",\n)\n" -+ } -+ }, -+ "crates__num-traits-0.2.19": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/num-traits/0.2.19/download" -+ ], -+ "strip_prefix": "num-traits-0.2.19", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"num_traits\",\n deps = [\n \"@crates__num-traits-0.2.19//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.2.19\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n deps = [\n \"@crates__autocfg-1.5.1//:autocfg\",\n ],\n edition = \"2021\",\n pkg_name = \"num-traits\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.2.19\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__paste-1.0.15": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/paste/1.0.15/download" -+ ], -+ "strip_prefix": "paste-1.0.15", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"paste\",\n deps = [\n \"@crates__paste-1.0.15//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.15\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"paste\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.15\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__proc-macro2-1.0.107": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/proc-macro2/1.0.107/download" -+ ], -+ "strip_prefix": "proc-macro2-1.0.107", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"proc_macro2\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:build_script_build\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.107\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"proc-macro2\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.107\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__quote-1.0.47": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/quote/1.0.47/download" -+ ], -+ "strip_prefix": "quote-1.0.47", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"quote\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.47\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"quote\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.47\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__regex-1.13.1": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/regex/1.13.1/download" -+ ], -+ "strip_prefix": "regex-1.13.1", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-automata-0.4.18//:regex_automata\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"perf\",\n \"perf-backtrack\",\n \"perf-cache\",\n \"perf-dfa\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-onepass\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.13.1\",\n)\n" -+ } -+ }, -+ "crates__regex-automata-0.4.18": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/regex-automata/0.4.18/download" -+ ], -+ "strip_prefix": "regex-automata-0.4.18", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_automata\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"dfa-onepass\",\n \"hybrid\",\n \"meta\",\n \"nfa-backtrack\",\n \"nfa-pikevm\",\n \"nfa-thompson\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-literal-multisubstring\",\n \"perf-literal-substring\",\n \"std\",\n \"syntax\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n \"unicode-word-boundary\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-automata\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.4.18\",\n)\n" -+ } -+ }, -+ "crates__regex-syntax-0.8.11": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/regex-syntax/0.8.11/download" -+ ], -+ "strip_prefix": "regex-syntax-0.8.11", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_syntax\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-syntax\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.8.11\",\n)\n" -+ } -+ }, -+ "crates__rustversion-1.0.23": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/rustversion/1.0.23/download" -+ ], -+ "strip_prefix": "rustversion-1.0.23", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"rustversion\",\n deps = [\n \"@crates__rustversion-1.0.23//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.23\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build/build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"rustversion\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.23\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -+ } -+ }, -+ "crates__syn-2.0.119": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/syn/2.0.119/download" -+ ], -+ "strip_prefix": "syn-2.0.119", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"extra-traits\",\n \"full\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.0.119\",\n)\n" -+ } -+ }, -+ "crates__syn-3.0.3": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/syn/3.0.3/download" -+ ], -+ "strip_prefix": "syn-3.0.3", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"3.0.3\",\n)\n" -+ } -+ }, -+ "crates__unicode-ident-1.0.24": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+ "patch_args": [], -+ "patch_tool": "", -+ "patches": [], -+ "remote_patch_strip": 1, -+ "sha256": "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75", -+ "type": "tar.gz", -+ "urls": [ -+ "https://static.crates.io/crates/unicode-ident/1.0.24/download" -+ ], -+ "strip_prefix": "unicode-ident-1.0.24", -+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"unicode_ident\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=unicode-ident\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.24\",\n)\n" -+ } -+ } -+ }, -+ "recordedRepoMappingEntries": [ -+ [ -+ "bazel_features+", -+ "bazel_features_globals", -+ "bazel_features++version_extension+bazel_features_globals" -+ ], -+ [ -+ "bazel_features+", -+ "bazel_features_version", -+ "bazel_features++version_extension+bazel_features_version" -+ ], -+ [ -+ "rules_cc+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "rules_cc+", -+ "cc_compatibility_proxy", -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -+ ], -+ [ -+ "rules_cc+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_features", -+ "bazel_features+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_skylib", -+ "bazel_skylib+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "rules_rust+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_rust+", -+ "rules_rust", -+ "rules_rust+" -+ ] -+ ] -+ } -+ }, -+ "@@rules_rust+//crate_universe/private:internal_extensions.bzl%cu_nr": { -+ "general": { -+ "bzlTransitiveDigest": "GOOgbXFJQhO4daGipwnspaixIHp6AWTvXRBe2wMULd4=", -+ "usagesDigest": "tG3p3Nb5XxC7vWY/bcKdb//g0HoAxpxxH3F5/jBVlk4=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+ "generatedRepoSpecs": { -+ "cargo_bazel_bootstrap": { -+ "repoRuleId": "@@rules_rust+//cargo/private:cargo_bootstrap.bzl%cargo_bootstrap_repository", -+ "attributes": { -+ "srcs": [ -+ "@@rules_rust+//crate_universe:src/api.rs", -+ "@@rules_rust+//crate_universe:src/api/lockfile.rs", -+ "@@rules_rust+//crate_universe:src/cli.rs", -+ "@@rules_rust+//crate_universe:src/cli/generate.rs", -+ "@@rules_rust+//crate_universe:src/cli/query.rs", -+ "@@rules_rust+//crate_universe:src/cli/render.rs", -+ "@@rules_rust+//crate_universe:src/cli/splice.rs", -+ "@@rules_rust+//crate_universe:src/cli/vendor.rs", -+ "@@rules_rust+//crate_universe:src/config.rs", -+ "@@rules_rust+//crate_universe:src/context.rs", -+ "@@rules_rust+//crate_universe:src/context/crate_context.rs", -+ "@@rules_rust+//crate_universe:src/context/platforms.rs", -+ "@@rules_rust+//crate_universe:src/lib.rs", -+ "@@rules_rust+//crate_universe:src/lockfile.rs", -+ "@@rules_rust+//crate_universe:src/main.rs", -+ "@@rules_rust+//crate_universe:src/metadata.rs", -+ "@@rules_rust+//crate_universe:src/metadata/cargo_bin.rs", -+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_resolver.rs", -+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.bat", -+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.sh", -+ "@@rules_rust+//crate_universe:src/metadata/dependency.rs", -+ "@@rules_rust+//crate_universe:src/metadata/metadata_annotation.rs", -+ "@@rules_rust+//crate_universe:src/rendering.rs", -+ "@@rules_rust+//crate_universe:src/rendering/template_engine.rs", -+ "@@rules_rust+//crate_universe:src/rendering/templates/module_bzl.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/header.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/aliases_map.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/deps_map.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_git.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_http.j2", -+ "@@rules_rust+//crate_universe:src/rendering/templates/vendor_module.j2", -+ "@@rules_rust+//crate_universe:src/rendering/verbatim/alias_rules.bzl", -+ "@@rules_rust+//crate_universe:src/select.rs", -+ "@@rules_rust+//crate_universe:src/splicing.rs", -+ "@@rules_rust+//crate_universe:src/splicing/cargo_config.rs", -+ "@@rules_rust+//crate_universe:src/splicing/crate_index_lookup.rs", -+ "@@rules_rust+//crate_universe:src/splicing/splicer.rs", -+ "@@rules_rust+//crate_universe:src/test.rs", -+ "@@rules_rust+//crate_universe:src/utils.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/glob.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/label.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/select.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/select_dict.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/select_list.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/select_scalar.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/select_set.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/serialize.rs", -+ "@@rules_rust+//crate_universe:src/utils/starlark/target_compatible_with.rs", -+ "@@rules_rust+//crate_universe:src/utils/symlink.rs", -+ "@@rules_rust+//crate_universe:src/utils/target_triple.rs" -+ ], -+ "binary": "cargo-bazel", -+ "cargo_lockfile": "@@rules_rust+//crate_universe:Cargo.lock", -+ "cargo_toml": "@@rules_rust+//crate_universe:Cargo.toml", -+ "version": "1.93.1", -+ "timeout": 900, -+ "rust_toolchain_cargo_template": "@rust_host_tools//:bin/{tool}", -+ "rust_toolchain_rustc_template": "@rust_host_tools//:bin/{tool}", -+ "compressed_windows_toolchain_names": false -+ } -+ } -+ }, -+ "moduleExtensionMetadata": { -+ "explicitRootModuleDirectDeps": [ -+ "cargo_bazel_bootstrap" -+ ], -+ "explicitRootModuleDirectDevDeps": [], -+ "useAllRepos": "NO", -+ "reproducible": false -+ }, -+ "recordedRepoMappingEntries": [ -+ [ -+ "bazel_features+", -+ "bazel_features_globals", -+ "bazel_features++version_extension+bazel_features_globals" -+ ], -+ [ -+ "bazel_features+", -+ "bazel_features_version", -+ "bazel_features++version_extension+bazel_features_version" -+ ], -+ [ -+ "rules_cc+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "rules_cc+", -+ "cc_compatibility_proxy", -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -+ ], -+ [ -+ "rules_cc+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_features", -+ "bazel_features+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_skylib", -+ "bazel_skylib+" -+ ], -+ [ -+ "rules_rust+", -+ "bazel_tools", -+ "bazel_tools" -+ ], -+ [ -+ "rules_rust+", -+ "cui", -+ "rules_rust++cu+cui" -+ ], -+ [ -+ "rules_rust+", -+ "rrc", -+ "rules_rust++i2+rrc" -+ ], -+ [ -+ "rules_rust+", -+ "rules_cc", -+ "rules_cc+" -+ ], -+ [ -+ "rules_rust+", -+ "rules_rust", -+ "rules_rust+" -+ ] -+ ] -+ } - } - }, - "facts": { diff --git a/third_party/xla/third_party/llvm/workspace.bzl b/third_party/xla/third_party/llvm/workspace.bzl index 7af3a8572d29aa..08ae305fe0537a 100644 --- a/third_party/xla/third_party/llvm/workspace.bzl +++ b/third_party/xla/third_party/llvm/workspace.bzl @@ -19,8 +19,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" - LLVM_COMMIT = "ab547095ead5464dc024d66264d9b8a987f429f3" - LLVM_SHA256 = "c47662ae375e9523870b87006f27c7f35f686f1489c0cbdf571a583a66c973c9" + LLVM_COMMIT = "cbc5a226cbf8d1d37d1ba8e55ce6973e8ef739cd" + LLVM_SHA256 = "1169d25c76701c535ecad40139ad48e4cec35d3e573308a2b5a2f63aaf2aa70a" tf_http_archive( name = name, diff --git a/third_party/xla/third_party/shardy/temporary.patch b/third_party/xla/third_party/shardy/temporary.patch index c49367187aedab..f9d67b31d32c2a 100644 --- a/third_party/xla/third_party/shardy/temporary.patch +++ b/third_party/xla/third_party/shardy/temporary.patch @@ -1,1560 +1,1499 @@ +diff --git a/third_party/llvm/build.patch b/third_party/llvm/build.patch +index e98a6a21..868226d5 100644 +--- a/third_party/llvm/build.patch ++++ b/third_party/llvm/build.patch +@@ -70,16 +70,3 @@ index a7e652c..5b8ac5e 100644 + "//conditions:default": [ + "BLAKE3_NO_AVX2", + "BLAKE3_NO_AVX512", +- +-diff --git a/utils/bazel/llvm-project-overlay/llvm/config.bzl b/utils/bazel/llvm-project-overlay/llvm/config.bzl +---- a/utils/bazel/llvm-project-overlay/llvm/config.bzl +-+++ b/utils/bazel/llvm-project-overlay/llvm/config.bzl +-@@ -110,6 +110,7 @@ +- # LLVM features +- r'LTDL_SHLIB_EXT=\".dll\"', +- r'LLVM_PLUGIN_EXT=\".dll\"', +-+ "LLVM_ENABLE_THREADS=1", +- ] + fenv_defines +- +- # TODO: We should switch to platforms-based config settings to make this easier +- diff --git a/third_party/llvm/generated.patch b/third_party/llvm/generated.patch -index f5b2c9f4..6258925a 100644 +index 6258925a..a490f07b 100644 --- a/third_party/llvm/generated.patch +++ b/third_party/llvm/generated.patch -@@ -1,18 +1,404 @@ --# Copyright 2026 The OpenXLA Authors. --# --# Licensed under the Apache License, Version 2.0 (the "License"); --# you may not use this file except in compliance with the License. --# You may obtain a copy of the License at --# --# http://www.apache.org/licenses/LICENSE-2.0 --# --# Unless required by applicable law or agreed to in writing, software --# distributed under the License is distributed on an "AS IS" BASIS, --# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. --# See the License for the specific language governing permissions and --# limitations under the License. --# ============================================================================== +@@ -1,65 +1,4 @@ Auto generated patch. Do not edit or delete it, even if empty. -+diff -ruN --strip-trailing-cr a/bolt/lib/Core/Relocation.cpp b/bolt/lib/Core/Relocation.cpp -+--- a/bolt/lib/Core/Relocation.cpp -++++ b/bolt/lib/Core/Relocation.cpp -+@@ -606,14 +606,6 @@ -+ case ELF::R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC: -+ case ELF::R_AARCH64_TLSLE_ADD_TPREL_HI12: -+ case ELF::R_AARCH64_TLSLE_ADD_TPREL_LO12_NC: -+- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12: -+- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12_NC: -+- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12: -+- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12_NC: -+- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12: -+- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC: -+- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12: -+- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC: -+ case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0: -+ case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0_NC: -+ case ELF::R_AARCH64_TLSDESC_LD64_LO12: -+diff -ruN --strip-trailing-cr a/bolt/test/AArch64/tls.c b/bolt/test/AArch64/tls.c -+--- a/bolt/test/AArch64/tls.c -++++ b/bolt/test/AArch64/tls.c -+@@ -5,8 +5,6 @@ -+ int b; -+ } tbssstruct = {}, tdatastruct = {4, 2}; -+ -+-__thread int directaccess; -+- -+ extern __thread struct str extstruct; -+ -+ extern void processAddr(volatile void *); -+@@ -20,9 +18,6 @@ -+ processAddr(&tbssstruct.b); -+ processAddr(&tdatastruct.b); -+ -+- // R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC for a direct access -+- directaccess++; -+- -+ // The R_AARCH64_TLSIE_ADR_GOTTPREL_PAGE21 and -+ // R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC relocations -+ processAddr(&extstruct.b); -+@@ -33,8 +28,6 @@ -+ // RUN: -Wl,--unresolved-symbols=ignore-all \ -+ // RUN: -fuse-ld=lld \ -+ // RUN: -nostdlib -+-// RUN: llvm-objdump -d -r --disassemble-symbols=main %t.exe \ -+-// RUN: | FileCheck %s --check-prefix=CHECK-DIRECT-ACCESS -+ // RUN: llvm-bolt %t.exe -o %t.bolt -+ // RUN: %clang %cflags -fPIC -pie %s -o %t_pie.exe -Wl,-q \ -+ // RUN: -Wl,--unresolved-symbols=ignore-all \ -+@@ -47,11 +40,6 @@ -+ // RUN: llvm-objdump -d -r --disassemble-symbols=main %t.so | FileCheck %s -+ // RUN: llvm-bolt %t.so -o %t.bolt.so -+ -+-// CHECK-DIRECT-ACCESS: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess -+-// CHECK-DIRECT-ACCESS-NEXT: add {{.*}} #0x1 -+-// CHECK-DIRECT-ACCESS-NEXT: str {{.*}} -+-// CHECK-DIRECT-ACCESS-NEXT: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess -+- -+ // Verify that unoptimized TLS access was generated for shared object. -+ // CHECK: adrp x0 -+ // CHECK-NEXT: R_AARCH64_TLSDESC_ADR_PAGE21 tbssstruct -+diff -ruN --strip-trailing-cr a/lldb/include/lldb/Symbol/Symbol.h b/lldb/include/lldb/Symbol/Symbol.h -+--- a/lldb/include/lldb/Symbol/Symbol.h -++++ b/lldb/include/lldb/Symbol/Symbol.h -+@@ -53,11 +53,7 @@ -+ -+ Symbol(const Symbol &rhs); -+ -+- ~Symbol() { -+- if (m_type != lldb::eSymbolTypeReExported && -+- m_type != lldb::eSymbolTypeInvalid) -+- m_addr_or_reexport.GetAddressRange(*this).Clear(); -+- } -++ ~Symbol() { m_addr_or_reexport.Clear(*this); } -+ -+ const Symbol &operator=(const Symbol &rhs); -+ -+@@ -419,6 +415,13 @@ -+ // impl to let the compiler know it's handled. -+ ~AddrRangeOrReExport() {} -+ -++ void Clear(const Symbol &sym) { -++ if (sym.GetType() == lldb::eSymbolTypeReExported) -++ m_reexport_info.Clear(); -++ else -++ m_addr_range.Clear(); -++ } -++ -+ private: -+ union { -+ // Contains the value, or the section offset address when the value is an -+diff -ruN --strip-trailing-cr a/lldb/source/API/SBThread.cpp b/lldb/source/API/SBThread.cpp -+--- a/lldb/source/API/SBThread.cpp -++++ b/lldb/source/API/SBThread.cpp -+@@ -465,7 +465,7 @@ -+ -+ // Release the run lock but keep the API lock. -+ TargetAPIMutex api_mutex = exe_ctx.AllowResume(); -+- std::lock_guard guard(api_mutex, std::adopt_lock); -++ std::unique_lock guard(api_mutex, std::adopt_lock); -+ if (process->GetTarget().GetDebugger().GetAsyncExecution()) -+ return process->Resume(); -+ return process->ResumeSynchronous(nullptr); -+diff -ruN --strip-trailing-cr a/lldb/unittests/Target/TargetAPIMutexTest.cpp b/lldb/unittests/Target/TargetAPIMutexTest.cpp -+--- a/lldb/unittests/Target/TargetAPIMutexTest.cpp -++++ b/lldb/unittests/Target/TargetAPIMutexTest.cpp -+@@ -120,6 +120,15 @@ -+ EXPECT_FALSE(background_lock.try_lock()); -+ }); -+ t.join(); -++ -++ // Unlock the original locked mutex. -++ // Calling try_lock() resolves the underlying mutex and re-enters it on this -++ // thread (incrementing the recursive count), so we unlock twice to fully -++ // release both acquisitions. -++ TargetAPIMutex cleanup_lock(target_sp); -++ ASSERT_TRUE(cleanup_lock.try_lock()); -++ cleanup_lock.unlock(); -++ cleanup_lock.unlock(); -+ } -+ -+ TEST_F(TargetAPIMutexTargetTest, LockGuardReleasesOnScopeExit) { -+diff -ruN --strip-trailing-cr a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp -+--- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp -++++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp -+@@ -5937,6 +5937,8 @@ -+ ISD::matchUnaryPredicate( -+ Y, -+ [&](auto *C) { -++ if (!C) -++ return true; -+ const APInt &YConst = C->getAsAPIntVal(); -+ return (Opcode == ISD::ABDS) -+ ? YConst.isSignedIntN(Bits) -+diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp -+--- a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp -++++ b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp -+@@ -1254,18 +1254,6 @@ -+ return true; -+ } -+ -+-/// Check whether \p GAN is the low part of a TLS address computation, i.e. the -+-/// second operand of an ADDlow. The target flags on their own do not tell the -+-/// ELF local-exec (:tprel_lo12: and :tprel_lo12_nc:) cases apart from other -+-/// uses, so callers that depend on local-exec semantics have to check the -+-/// object format as well. Local dynamic never gets here because it does not -+-/// build an ADDlow. -+-static bool isTLSLo12(const GlobalAddressSDNode *GAN) { -+- unsigned Flags = GAN->getTargetFlags(); -+- return (Flags & (AArch64II::MO_TLS | AArch64II::MO_FRAGMENT)) == -+- (AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); -+-} -+- -+ /// Check if the immediate offset is valid as a scaled immediate. -+ static bool isValidAsScaledImmediate(int64_t Offset, unsigned Range, -+ unsigned Size) { -+@@ -1361,13 +1349,8 @@ -+ if (!GAN) -+ return true; -+ -+- // Folding the low part of an ELF local-exec TLS address into a 128-bit -+- // access needs R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC variant, which -+- // the GNU bfd linker does not support, so keep materialising the address -+- // with an add. -+ if (GAN->getOffset() % Size == 0 && -+- GAN->getGlobal()->getPointerAlignment(DL) >= Size && -+- !(Size > 8 && Subtarget->isTargetELF() && isTLSLo12(GAN))) -++ GAN->getGlobal()->getPointerAlignment(DL) >= Size) -+ return true; -+ } -+ -+diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp -+--- a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp -++++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp -+@@ -11542,7 +11542,10 @@ -+ // add x0, x0, :tprel_lo12:a -+ SDValue Var = DAG.getTargetGlobalAddress( -+ GV, DL, PtrVT, 0, AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); -+- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, ThreadBase, Var); -++ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, ThreadBase, -++ Var, -++ DAG.getTargetConstant(0, DL, MVT::i32)), -++ 0); -+ } -+ -+ case 24: { -+@@ -11558,10 +11561,9 @@ -+ HiVar, -+ DAG.getTargetConstant(0, DL, MVT::i32)), -+ 0); -+- // Emit the low part as an ADDlow so that it can be folded into the -+- // addressing mode of a following load or store, turning the add into a -+- // :tprel_lo12_nc: relocation on the memory access itself. -+- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, Addr, LoVar); -++ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, Addr, LoVar, -++ DAG.getTargetConstant(0, DL, MVT::i32)), -++ 0); -+ } -+ -+ case 32: { -+diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/abd-combine.ll b/llvm/test/CodeGen/AArch64/abd-combine.ll -+--- a/llvm/test/CodeGen/AArch64/abd-combine.ll -++++ b/llvm/test/CodeGen/AArch64/abd-combine.ll -+@@ -523,9 +523,26 @@ -+ ret <1 x i64> %10 -+ } -+ -++; Poison elements in the constant operand pass a null ConstantSDNode to the -++; matchUnaryPredicate lambda in visitABD; make sure that does not crash. -++define <4 x i32> @abdu_const_poison(<4 x i8> %x) { -++; CHECK-LABEL: abdu_const_poison: -++; CHECK: // %bb.0: -++; CHECK-NEXT: mov w8, #8388736 // =0x800080 -++; CHECK-NEXT: bic v0.4h, #255, lsl #8 -++; CHECK-NEXT: fmov d1, x8 -++; CHECK-NEXT: uabdl v0.4s, v0.4h, v1.4h -++; CHECK-NEXT: ret -++ %ext = zext <4 x i8> %x to <4 x i32> -++ %sub = sub <4 x i32> , %ext -++ %abs = call <4 x i32> @llvm.abs.v4i32(<4 x i32> %sub, i1 false) -++ ret <4 x i32> %abs -++} -++ -+ declare <8 x i8> @llvm.aarch64.neon.umax.v8i8(<8 x i8>, <8 x i8>) -+ declare <1 x i64> @llvm.aarch64.neon.saddlp.v1i64.v2i32(<2 x i32>) -+ declare <8 x i8> @llvm.aarch64.neon.uabd.v8i8(<8 x i8>, <8 x i8>) -+ declare <8 x i16> @llvm.aarch64.neon.uabd.v8i16(<8 x i16>, <8 x i16>) -+ declare <8 x i16> @llvm.aarch64.neon.sabd.v8i16(<8 x i16>, <8 x i16>) -+ declare <8 x i32> @llvm.abs.v8i32(<8 x i32>, i1) -++declare <4 x i32> @llvm.abs.v4i32(<4 x i32>, i1) -+diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll -+--- a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll -++++ b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll -+@@ -27,24 +27,24 @@ -+ ; RUN: llc -mtriple=arm64-none-linux-gnu -filetype=obj < %s -code-model=large | llvm-objdump -r - | FileCheck --check-prefix=CHECK-24-RELOC %s -+ -+ @local_exec_var = thread_local(localexec) global i32 0 -+-@local_exec_var64 = thread_local(localexec) global i64 0 -+-@vec_local_exec_var = thread_local(localexec) global <2 x i64> zeroinitializer, align 16 -+ -+ define i32 @test_local_exec() { -+ ; CHECK-LABEL: test_local_exec: -+ %val = load i32, ptr @local_exec_var -+ -+ ; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+-; CHECK-12: ldr w0, [x[[R1]], :tprel_lo12:local_exec_var] -++; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:local_exec_var -++; CHECK-12: ldr w0, [x[[R2]]] -+ -+-; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12 -++; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 -+ -+ ; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+ ; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var -+-; CHECK-24: ldr w0, [x[[R2]], :tprel_lo12_nc:local_exec_var] -++; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:local_exec_var -++; CHECK-24: ldr w0, [x[[R3]]] -+ -+ ; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 -+-; CHECK-24-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC -++; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC -+ -+ ; CHECK-32: movz x[[R2:[0-9]+]], #:tprel_g1:local_exec_var -+ ; CHECK-32: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+@@ -66,24 +66,6 @@ -+ ret i32 %val -+ } -+ -+-define void @test_local_exec_store64(i64 %val) { -+-; CHECK-LABEL: test_local_exec_store64: -+- store i64 %val, ptr @local_exec_var64 -+- -+-; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+-; CHECK-12: str x0, [x[[R1]], :tprel_lo12:local_exec_var64] -+- -+-; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST64_TPREL_LO12 -+- -+-; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+-; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var64 -+-; CHECK-24: str x0, [x[[R2]], :tprel_lo12_nc:local_exec_var64] -+- -+-; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 local_exec_var64 -+-; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC local_exec_var64 -+- ret void -+-} -+- -+ define ptr @test_local_exec_addr() { -+ ; CHECK-LABEL: test_local_exec_addr: -+ ret ptr @local_exec_var -+@@ -122,26 +104,3 @@ -+ ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G1_NC -+ ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G0_NC -+ } -+- -+-; A 128-bit access would need R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC -+-; variant, which not every linker implements, so the low part stays in a -+-; separate add. -+-define <2 x i64> @test_local_exec_128bit() { -+-; CHECK-LABEL: test_local_exec_128bit: -+- %val = load <2 x i64>, ptr @vec_local_exec_var -+- -+-; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+-; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:vec_local_exec_var -+-; CHECK-12: ldr q0, [x[[R2]]] -+- -+-; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 vec_local_exec_var -+- -+-; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 -+-; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:vec_local_exec_var -+-; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:vec_local_exec_var -+-; CHECK-24: ldr q0, [x[[R3]]] -+- -+-; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 vec_local_exec_var -+-; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC vec_local_exec_var -+- ret <2 x i64> %val -+-} -+diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/win-tls.ll b/llvm/test/CodeGen/AArch64/win-tls.ll -+--- a/llvm/test/CodeGen/AArch64/win-tls.ll -++++ b/llvm/test/CodeGen/AArch64/win-tls.ll -+@@ -3,7 +3,6 @@ -+ @tlsVar = thread_local global i32 0 -+ @tlsVar8 = thread_local global i8 0 -+ @tlsVar64 = thread_local global i64 0 -+-@tlsVar128 = thread_local global <2 x i64> zeroinitializer -+ -+ define i32 @getVar() { -+ %1 = load i32, ptr @tlsVar -+@@ -29,11 +28,6 @@ -+ ret i64 %1 -+ } -+ -+-define <2 x i64> @getVar128() { -+- %1 = load <2 x i64>, ptr @tlsVar128 -+- ret <2 x i64> %1 -+-} -+- -+ ; CHECK-LABEL: getVar -+ ; CHECK: adrp [[TLS_INDEX_ADDR:x[0-9]+]], _tls_index -+ ; CHECK: ldr [[TLS_POINTER:x[0-9]+]], [x18, #88] -+@@ -68,7 +62,3 @@ -+ ; CHECK-LABEL: getVar64 -+ ; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar64 -+ ; CHECK: ldr x0, [[[TLS]], :secrel_lo12:tlsVar64] -+- -+-; CHECK-LABEL: getVar128 -+-; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar128 -+-; CHECK: ldr q0, [[[TLS]], :secrel_lo12:tlsVar128] -+diff -ruN --strip-trailing-cr a/utils/bazel/.bazelrc b/utils/bazel/.bazelrc -+--- a/utils/bazel/.bazelrc -++++ b/utils/bazel/.bazelrc -+@@ -222,6 +222,19 @@ -+ build:hermetic-toolchain --copt=-Wno-modules-import-nested-redundant --host_copt=-Wno-modules-import-nested-redundant -+ -+ ############################################################################### -++# Options for Emscripten WebAssembly builds. -++############################################################################### -++ -++build:wasm --platforms=@emsdk//:platform_wasm -++# Match LLVM's single-threaded config and shut the runtime down when main returns. -++build:wasm --features=-use_pthreads,exit_runtime -++# Make the default CLI artifact use Node's host filesystem; browser builds override this to 0. -++build:wasm --linkopt=-sNODERAWFS=1 -++ -++# TODO: zstd unconditionally enables pthreads on non-Windows targets. -++build:wasm --@llvm-project//third-party:llvm_enable_zstd=false -++ -++############################################################################### -+ # Options for continuous integration. -+ ############################################################################### -+ -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel -+--- a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel -++++ b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel -+@@ -1957,6 +1957,7 @@ -+ ":parse", -+ ":sema", -+ ":serialization", -++ "//compiler-rt:emutls", -+ "//llvm:AllTargetsAsmParsers", -+ "//llvm:AllTargetsCodeGens", -+ "//llvm:Core", -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel -+--- a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel -++++ b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel -+@@ -165,6 +165,15 @@ -+ ], -+ ) -+ -++cc_library( -++ name = "emutls", -++ srcs = [ -++ "lib/builtins/emutls.c", -++ ], -++ hdrs = glob(["lib/builtins/*.h"]), -++ linkstatic = True, -++) -++ -+ filegroup( -+ name = "fuzzer_installed_hdrs", -+ srcs = glob(["include/fuzzer/*.h"]), - diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel - --- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel - +++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel -@@ -35,23 +421,1084 @@ diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.baze - diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl - --- a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl - +++ b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl --@@ -61,14 +61,19 @@ -- else: -- deps = deps + ["//libc/test/UnitTest:LibcUnitTest"] +-diff -ruN --strip-trailing-cr a/bolt/lib/Core/Relocation.cpp b/bolt/lib/Core/Relocation.cpp +---- a/bolt/lib/Core/Relocation.cpp +-+++ b/bolt/lib/Core/Relocation.cpp +-@@ -606,14 +606,6 @@ +- case ELF::R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC: +- case ELF::R_AARCH64_TLSLE_ADD_TPREL_HI12: +- case ELF::R_AARCH64_TLSLE_ADD_TPREL_LO12_NC: +-- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12: +-- case ELF::R_AARCH64_TLSLE_LDST8_TPREL_LO12_NC: +-- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12: +-- case ELF::R_AARCH64_TLSLE_LDST16_TPREL_LO12_NC: +-- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12: +-- case ELF::R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC: +-- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12: +-- case ELF::R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC: +- case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0: +- case ELF::R_AARCH64_TLSLE_MOVW_TPREL_G0_NC: +- case ELF::R_AARCH64_TLSDESC_LD64_LO12: +-diff -ruN --strip-trailing-cr a/bolt/test/AArch64/tls.c b/bolt/test/AArch64/tls.c +---- a/bolt/test/AArch64/tls.c +-+++ b/bolt/test/AArch64/tls.c +-@@ -5,8 +5,6 @@ +- int b; +- } tbssstruct = {}, tdatastruct = {4, 2}; +- +--__thread int directaccess; +-- +- extern __thread struct str extstruct; +- +- extern void processAddr(volatile void *); +-@@ -20,9 +18,6 @@ +- processAddr(&tbssstruct.b); +- processAddr(&tdatastruct.b); - --+ tags = kwargs.pop("tags", []) -- if full_build: -+@@ -66,7 +66,7 @@ - copts = copts + _FULL_BUILD_COPTS +-- // R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC for a direct access +-- directaccess++; +-- +- // The R_AARCH64_TLSIE_ADR_GOTTPREL_PAGE21 and +- // R_AARCH64_TLSIE_LD64_GOTTPREL_LO12_NC relocations +- processAddr(&extstruct.b); +-@@ -33,8 +28,6 @@ +- // RUN: -Wl,--unresolved-symbols=ignore-all \ +- // RUN: -fuse-ld=lld \ +- // RUN: -nostdlib +--// RUN: llvm-objdump -d -r --disassemble-symbols=main %t.exe \ +--// RUN: | FileCheck %s --check-prefix=CHECK-DIRECT-ACCESS +- // RUN: llvm-bolt %t.exe -o %t.bolt +- // RUN: %clang %cflags -fPIC -pie %s -o %t_pie.exe -Wl,-q \ +- // RUN: -Wl,--unresolved-symbols=ignore-all \ +-@@ -47,11 +40,6 @@ +- // RUN: llvm-objdump -d -r --disassemble-symbols=main %t.so | FileCheck %s +- // RUN: llvm-bolt %t.so -o %t.bolt.so +- +--// CHECK-DIRECT-ACCESS: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess +--// CHECK-DIRECT-ACCESS-NEXT: add {{.*}} #0x1 +--// CHECK-DIRECT-ACCESS-NEXT: str {{.*}} +--// CHECK-DIRECT-ACCESS-NEXT: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC directaccess +-- +- // Verify that unoptimized TLS access was generated for shared object. +- // CHECK: adrp x0 +- // CHECK-NEXT: R_AARCH64_TLSDESC_ADR_PAGE21 tbssstruct + diff -ruN --strip-trailing-cr a/lldb/include/lldb/Symbol/Symbol.h b/lldb/include/lldb/Symbol/Symbol.h + --- a/lldb/include/lldb/Symbol/Symbol.h + +++ b/lldb/include/lldb/Symbol/Symbol.h +@@ -90,41 +29,10 @@ diff -ruN --strip-trailing-cr a/lldb/include/lldb/Symbol/Symbol.h b/lldb/include + private: + union { + // Contains the value, or the section offset address when the value is an +-diff -ruN --strip-trailing-cr a/lldb/source/API/SBThread.cpp b/lldb/source/API/SBThread.cpp +---- a/lldb/source/API/SBThread.cpp +-+++ b/lldb/source/API/SBThread.cpp +-@@ -465,7 +465,7 @@ +- +- // Release the run lock but keep the API lock. +- TargetAPIMutex api_mutex = exe_ctx.AllowResume(); +-- std::lock_guard guard(api_mutex, std::adopt_lock); +-+ std::unique_lock guard(api_mutex, std::adopt_lock); +- if (process->GetTarget().GetDebugger().GetAsyncExecution()) +- return process->Resume(); +- return process->ResumeSynchronous(nullptr); +-diff -ruN --strip-trailing-cr a/lldb/unittests/Target/TargetAPIMutexTest.cpp b/lldb/unittests/Target/TargetAPIMutexTest.cpp +---- a/lldb/unittests/Target/TargetAPIMutexTest.cpp +-+++ b/lldb/unittests/Target/TargetAPIMutexTest.cpp +-@@ -120,6 +120,15 @@ +- EXPECT_FALSE(background_lock.try_lock()); +- }); +- t.join(); -+ --+ # Temporarily disable full_build tests (currently broken) to unblock CI. --+ tags = tags + ["manual", "notap"] -+ -+ # Temporarily disable full_build tests (currently broken) to unblock CI. -+- tags = tags + ["manual", "notap"] -++ tags = tags + ["manual", "nobuildkite", "notap"] - cc_test( - name = name, - local_defines = local_defines + _TEST_DEFINES + LIBC_CONFIGURE_OPTIONS, -- deps = deps, -- copts = copts + libc_common_copts(), -- linkstatic = 1, --+ tags = tags, -- **kwargs -- ) -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel -+--- a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel -++++ b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel -+@@ -456,6 +456,7 @@ -+ "-pthread", -+ "-ldl", -+ ], -++ "@platforms//os:emscripten": [], -+ "//conditions:default": [ -+ "-pthread", -+ "-ldl", -+@@ -1439,7 +1440,10 @@ -+ "include/llvm/Analysis/Utils/*.h", -+ ], -+ ) + ["include/llvm-c/Analysis.h"], -+- copts = llvm_copts + ["-ftrapping-math"], -++ copts = llvm_copts + select({ -++ "@platforms//os:emscripten": [], -++ "//conditions:default": ["-ftrapping-math"], -++ }), -+ features = ["-parse_headers"], -+ textual_hdrs = glob([ -+ "include/llvm/Analysis/*.def", -+@@ -4973,6 +4977,7 @@ -+ # ll scripts rely on symbols from dependent -+ # libraries being resolvable. -+ linkopts = select({ -++ "@platforms//os:emscripten": [], -+ "@platforms//os:macos": [], -+ "@platforms//os:windows": [], -+ "//conditions:default": [ -+@@ -5548,6 +5553,7 @@ -+ copts = llvm_copts, -+ # Make symbols from the standard library dynamically resolvable. -+ linkopts = select({ -++ "@platforms//os:emscripten": [], -+ "@platforms//os:macos": [], -+ "@platforms//os:windows": [], -+ "//conditions:default": [ -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/config.bzl b/utils/bazel/llvm-project-overlay/llvm/config.bzl -+--- a/utils/bazel/llvm-project-overlay/llvm/config.bzl -++++ b/utils/bazel/llvm-project-overlay/llvm/config.bzl -+@@ -46,7 +46,28 @@ -+ "HAVE_UNISTD_H=1", -+ ] -+ -++emscripten_defines = [ -++ "LLVM_ON_UNIX=1", -++ r'LTDL_SHLIB_EXT=\".so\"', -++ r'LLVM_PLUGIN_EXT=\".so\"', -++ "LLVM_ENABLE_LLVM_EXPORT_ANNOTATIONS=1", -++ "LLVM_ENABLE_PLUGINS=0", -++ "LLVM_ENABLE_THREADS=0", -++ "HAVE_MALLINFO=1", -++ "HAVE_SETENV_R=1", -++ "HAVE_STRERROR_R=1", -++ "HAVE_SYSEXITS_H=1", -++ "HAVE_SYS_IOCTL_H=1", -++ "HAVE_UNISTD_H=1", -++] -++ -++fenv_defines = [ -++ "HAVE_DECL_FE_ALL_EXCEPT=1", -++ "HAVE_DECL_FE_INEXACT=1", -++] -++ -+ backtrace_defines = select({ -++ "@platforms//os:emscripten": [], -+ "@platforms//os:windows": [], -+ "@llvm//platforms/config:musl": [], -+ "//conditions:default": [ -+@@ -60,14 +81,14 @@ -+ "//conditions:default": [], -+ }) -+ -+-linux_defines = posix_defines + [ -++linux_defines = posix_defines + fenv_defines + [ -+ "_GNU_SOURCE", -+ "HAVE_GETAUXVAL=1", -+ "HAVE_SBRK=1", -+ "HAVE_STRUCT_STAT_ST_MTIM_TV_NSEC=1", -+ ] -+ -+-macos_defines = posix_defines + [ -++macos_defines = posix_defines + fenv_defines + [ -+ "HAVE_MACH_MACH_H=1", -+ "HAVE_MALLOC_MALLOC_H=1", -+ "HAVE_MALLOC_ZONE_STATISTICS=1", -+@@ -89,12 +110,13 @@ -+ # LLVM features -+ r'LTDL_SHLIB_EXT=\".dll\"', -+ r'LLVM_PLUGIN_EXT=\".dll\"', -+-] -++] + fenv_defines -+ -+ # TODO: We should switch to platforms-based config settings to make this easier -+ # to express. -+ os_defines = select({ -+- "@platforms//os:freebsd": posix_defines, -++ "@platforms//os:emscripten": emscripten_defines, -++ "@platforms//os:freebsd": posix_defines + fenv_defines, -+ "@platforms//os:macos": macos_defines, -+ "@platforms//os:windows": win32_defines, -+ "//conditions:default": linux_defines, -+@@ -117,6 +139,7 @@ -+ Label("//llvm:linux_ppc64le"): native_arch_defines("PowerPC", "powerpc64le-unknown-linux-gnu"), -+ Label("//llvm:linux_riscv64"): native_arch_defines("RISCV", "riscv64-unknown-linux-gnu"), -+ Label("//llvm:linux_s390x"): native_arch_defines("SystemZ", "systemz-unknown-linux_gnu"), -++ "@platforms//os:emscripten": native_arch_defines("WebAssembly", "wasm32-unknown-emscripten"), -+ "@platforms//os:windows": native_arch_defines("X86", "x86_64-pc-win32"), -+ "//conditions:default": native_arch_defines("X86", "x86_64-unknown-linux-gnu"), -+ }) + [ -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h -+--- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h -++++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h -+@@ -56,11 +56,11 @@ -+ -+ /* Define to 1 if you have the declaration of `FE_ALL_EXCEPT', and to 0 if you -+ don't. */ -+-#define HAVE_DECL_FE_ALL_EXCEPT 1 -++/* HAVE_DECL_FE_ALL_EXCEPT defined in Bazel */ -+ -+ /* Define to 1 if you have the declaration of `FE_INEXACT', and to 0 if you -+ don't. */ -+-#define HAVE_DECL_FE_INEXACT 1 -++/* HAVE_DECL_FE_INEXACT defined in Bazel */ -+ -+ /* Define to 1 if you have the declaration of `strerror_s', and to 0 if you -+ don't. */ -+diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h -+--- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h -++++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h -+@@ -28,7 +28,7 @@ -+ /* LLVM_DEFAULT_TARGET_TRIPLE defined in Bazel */ -+ -+ /* Define if threads enabled */ -+-#define LLVM_ENABLE_THREADS 1 -++/* LLVM_ENABLE_THREADS defined in Bazel */ -+ -+ /* Has gcc/MSVC atomic intrinsics */ -+ #define LLVM_HAS_ATOMICS 1 -+diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel b/utils/bazel/MODULE.bazel -+--- a/utils/bazel/MODULE.bazel -++++ b/utils/bazel/MODULE.bazel -+@@ -30,6 +30,7 @@ -+ bazel_dep(name = "libpfm", version = "4.13.0", repo_name = "pfm") -+ bazel_dep(name = "vulkan_headers", version = "1.4.349") -+ -++bazel_dep(name = "emsdk", version = "6.0.2", dev_dependency = True) -+ bazel_dep(name = "llvm", version = "0.8.5", dev_dependency = True) - -+ llvm_repos_extension = use_extension(":extensions.bzl", "llvm_repos_extension") -+diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel.lock b/utils/bazel/MODULE.bazel.lock -+--- a/utils/bazel/MODULE.bazel.lock -++++ b/utils/bazel/MODULE.bazel.lock -+@@ -81,6 +81,8 @@ -+ "https://bcr.bazel.build/modules/buildozer/7.1.2/source.json": "c9028a501d2db85793a6996205c8de120944f50a0d570438fcae0457a5f9d1f8", -+ "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/MODULE.bazel": "f6561baff0fc0035c9c1a9e2b0820de106cdb01b37bf5c81276860ccc863e5b2", -+ "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/source.json": "a8611a2b5577929ad7e1f44ded19dab21a188125a74ac6192d21d283609f280f", -++ "https://bcr.bazel.build/modules/emsdk/6.0.2/MODULE.bazel": "4a3c4195e5f2e0056bc18bf9f8af631c4f720c3e8cea45cfb7247eacf02e27fe", -++ "https://bcr.bazel.build/modules/emsdk/6.0.2/source.json": "53111cbcb9f0971aa14da976bafbb937a1bc92a2580c8881890983cd2d289af0", -+ "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/MODULE.bazel": "f1b7bb2dd53e8f2ef984b39485ec8a44e9076dda5c4b8efd2fb4c6a6e856a31d", -+ "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/source.json": "ebe931bfe362e4b41e59ee00a528db6074157ff2ced92eb9e970acab2e1089c9", -+ "https://bcr.bazel.build/modules/gazelle/0.32.0/MODULE.bazel": "b499f58a5d0d3537f3cf5b76d8ada18242f64ec474d8391247438bf04f58c7b8", -+@@ -184,6 +186,7 @@ -+ "https://bcr.bazel.build/modules/rules_cc/0.2.13/MODULE.bazel": "eecdd666eda6be16a8d9dc15e44b5c75133405e820f620a234acc4b1fdc5aa37", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.14/MODULE.bazel": "353c99ed148887ee89c54a17d4100ae7e7e436593d104b668476019023b58df8", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.15/MODULE.bazel": "6a0a4a75a57aa6dc888300d848053a58c6b12a29f89d4304e1c41448514ec6e8", -++ "https://bcr.bazel.build/modules/rules_cc/0.2.16/MODULE.bazel": "9242fa89f950c6ef7702801ab53922e99c69b02310c39fb6e62b2bd30df2a1d4", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.17/MODULE.bazel": "1849602c86cb60da8613d2de887f9566a6d354a6df6d7009f9d04a14402f9a84", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.18/MODULE.bazel": "4460ec36adc8f722a6a2a4ac9374cb91f2acebadaa93fc37966129afb3dece87", -+ "https://bcr.bazel.build/modules/rules_cc/0.2.19/MODULE.bazel": "d5e0f05b63273281a16654eb6b1a8742a75ec153ac8b4f0419949d6e401e46f0", -+@@ -242,6 +245,8 @@ -+ "https://bcr.bazel.build/modules/rules_license/0.0.7/MODULE.bazel": "088fbeb0b6a419005b89cf93fe62d9517c0a2b8bb56af3244af65ecfe37e7d5d", -+ "https://bcr.bazel.build/modules/rules_license/1.0.0/MODULE.bazel": "a7fda60eefdf3d8c827262ba499957e4df06f659330bbe6cdbdb975b768bb65c", -+ "https://bcr.bazel.build/modules/rules_license/1.0.0/source.json": "a52c89e54cc311196e478f8382df91c15f7a2bfdf4c6cd0e2675cc2ff0b56efb", -++ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/MODULE.bazel": "c22a48b2a0dbf05a9dc5f83837bbc24c226c1f6e618de3c3a610044c9f336056", -++ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/source.json": "a3f966f4415a8a6545e560ee5449eac95cc633f96429d08e87c87775c72f5e09", -+ "https://bcr.bazel.build/modules/rules_pkg/0.7.0/MODULE.bazel": "df99f03fc7934a4737122518bb87e667e62d780b610910f0447665a7e2be62dc", -+ "https://bcr.bazel.build/modules/rules_pkg/1.0.1/MODULE.bazel": "5b1df97dbc29623bccdf2b0dcd0f5cb08e2f2c9050aab1092fd39a41e82686ff", -+ "https://bcr.bazel.build/modules/rules_pkg/1.0.1/source.json": "bd82e5d7b9ce2d31e380dd9f50c111d678c3bdaca190cb76b0e1c71b05e1ba8a", -+@@ -268,7 +273,8 @@ -+ "https://bcr.bazel.build/modules/rules_python/1.6.0/MODULE.bazel": "7e04ad8f8d5bea40451cf80b1bd8262552aa73f841415d20db96b7241bd027d8", -+ "https://bcr.bazel.build/modules/rules_python/1.7.0/MODULE.bazel": "d01f995ecd137abf30238ad9ce97f8fc3ac57289c8b24bd0bf53324d937a14f8", -+ "https://bcr.bazel.build/modules/rules_python/1.8.0/MODULE.bazel": "c151c025dbcc93d8f62ab68ecc313c9176a868a0e6386981bf2a12aec77cbe7b", -+- "https://bcr.bazel.build/modules/rules_python/1.8.0/source.json": "356397eed5b46971d8c585c92098d70495078a80bf18bebcb4209f44b495f3e6", -++ "https://bcr.bazel.build/modules/rules_python/1.8.4/MODULE.bazel": "33e3971e66161a3e955f7a0d411a8d1f291c4ce4c561851512466f3c77ff8ece", -++ "https://bcr.bazel.build/modules/rules_python/1.8.4/source.json": "9fbc0e57bae52cddcc3831d668bce87a47e0c655104a85098d4459dd9a3b0a10", -+ "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/MODULE.bazel": "d44fec647d0aeb67b9f3b980cf68ba634976f3ae7ccd6c07d790b59b87a4f251", -+ "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/source.json": "37c10335f2361c337c5c1f34ed36d2da70534c23088062b33a8bdaab68aa9dea", -+ "https://bcr.bazel.build/modules/rules_rust/0.69.0/MODULE.bazel": "4326fec48f2fef0d514de46346f7f77e200c82936dd08b91c9ef039fbdad5c10", -+@@ -348,9 +354,142 @@ -+ ] -+ } -+ }, -++ "@@emsdk+//:emscripten_cache.bzl%emscripten_cache": { -++ "general": { -++ "bzlTransitiveDigest": "GMscy7c4sDbvbf9dMSrtUvyJJBuxYpJBT/Lg+2ob6dk=", -++ "usagesDigest": "Id/C4z1d3MUlKZgmBOQiLi2E7NQZTGsd4WV2wSMzmo4=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": {}, -++ "generatedRepoSpecs": { -++ "emscripten_cache": { -++ "repoRuleId": "@@emsdk+//:emscripten_cache.bzl%_emscripten_cache_repository", -++ "attributes": { -++ "configuration": [], -++ "targets": [], -++ "prebuilt_cache_url": "", -++ "prebuilt_cache_sha256": "", -++ "prebuilt_cache_strip_prefix": "" -++ } -++ } -++ }, -++ "recordedRepoMappingEntries": [] -++ } -++ }, -++ "@@emsdk+//:emscripten_deps.bzl%emscripten_deps": { -++ "general": { -++ "bzlTransitiveDigest": "ZT33Pf8H8gJ/X4gT3oXVPzuAO0H2y1HtHMPOfILGV0M=", -++ "usagesDigest": "4SlUap0Npa9PDUrLoi0uZ7CRDp9h/YbWi0UuidttuWc=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": {}, -++ "generatedRepoSpecs": { -++ "emscripten_bin_linux": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -++ "sha256": "d574428df9ecf00790e28636bdc47027432737c31621b18cdb418123afda4ac1", -++ "strip_prefix": "install", -++ "type": "tar.xz", -++ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" -++ } -++ }, -++ "emscripten_bin_linux_arm64": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -++ "sha256": "d74803ef563511b9cc1e5cde5016f06d161ffd2b6223135a8aeeef44194594e7", -++ "strip_prefix": "install", -++ "type": "tar.xz", -++ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" -++ } -++ }, -++ "emscripten_bin_mac": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -++ "sha256": "356f36ba04a54edb029c658dd1b547c5c8a8f3c166b09654c1efcb9cf7bf8a57", -++ "strip_prefix": "install", -++ "type": "tar.xz", -++ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" -++ } -++ }, -++ "emscripten_bin_mac_arm64": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -++ "sha256": "ded3bb783e7aa3dda576955dd0aa3a71dd21789e42befb63ee14f7d9f9b6aa32", -++ "strip_prefix": "install", -++ "type": "tar.xz", -++ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" -++ } -++ }, -++ "emscripten_bin_win": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/clang++.exe\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/llvm-ar.exe\",\n \"bin/llvm-dwarfdump.exe\",\n \"bin/llvm-nm.exe\",\n \"bin/llvm-objcopy.exe\",\n \"bin/wasm-ctor-eval.exe\",\n \"bin/wasm-emscripten-finalize.exe\",\n \"bin/wasm-ld.exe\",\n \"bin/wasm-metadce.exe\",\n \"bin/wasm-opt.exe\",\n \"bin/wasm-split.exe\",\n \"bin/wasm2js.exe\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar.exe\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", -++ "sha256": "e5f9250a9cf4ff6ed16d57d6b5e177c844067381d31cd0c1a607c1ee1d2ba088", -++ "strip_prefix": "install", -++ "type": "zip", -++ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/win/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.zip" -++ } -++ } -++ }, -++ "recordedRepoMappingEntries": [ -++ [ -++ "emsdk+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "emsdk+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_cc+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "rules_cc+", -++ "cc_compatibility_proxy", -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -++ ], -++ [ -++ "rules_cc+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -++ "rules_cc", -++ "rules_cc+" -++ ] -++ ] -++ } -++ }, -++ "@@protobuf+//python/dist:system_python.bzl%system_python_extension": { -++ "general": { -++ "bzlTransitiveDigest": "qh0n9IrXU/xS94wxKQrG1J63zrLkA1Wy2Y3BQxptPcI=", -++ "usagesDigest": "tCi55FyqtOJ2jXh9vcjrHCl4ov3kpWiwKl103nA9BOI=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": {}, -++ "generatedRepoSpecs": { -++ "system_python": { -++ "repoRuleId": "@@protobuf+//python/dist:system_python.bzl%system_python", -++ "attributes": { -++ "minimum_python_version": "3.9" -++ } -++ } -++ }, -++ "recordedRepoMappingEntries": [] -++ } -++ }, -+ "@@pybind11_bazel+//:internal_configure.bzl%internal_configure_extension": { -+ "general": { -+- "bzlTransitiveDigest": "NFQjcZF+fAvf5fDH+pqsx4JrfzP9PuHBz6S6ZutIbnw=", -++ "bzlTransitiveDigest": "7zBsfo5dyMqKT23rXrvWqJMx0AugwL6NyirkmvzKcqU=", -+ "usagesDigest": "D1r3lfzMuUBFxgG8V6o0bQTLMk3GkaGOaPzw53wrwyw=", -+ "recordedFileInputs": { -+ "@@pybind11_bazel+//MODULE.bazel": "e6f4c20442eaa7c90d7190d8dc539d0ab422f95c65a57cc59562170c58ae3d34" -+@@ -380,8 +519,8 @@ -+ }, -+ "@@rules_android+//bzlmod_extensions:apksig.bzl%apksig_extension": { -+ "general": { -+- "bzlTransitiveDigest": "By9qVNN7G4oL1vYOJXye7Dp/CbR2ar9oxAW8WXAVcVw=", -+- "usagesDigest": "xq6OVkELeJvOgYo3oY/sUBsGFbcqdV+9BYiNgSPV/po=", -++ "bzlTransitiveDigest": "15xx/lo4VYL9KdLW0Cc94ebALMF07+XZH8dZcVU8/LI=", -++ "usagesDigest": "S8lLnnZxdeYUYq3kIGhVMk0wQ9Fd6elmCskvn+SL6iw=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+@@ -389,7 +528,10 @@ -+ "apksig": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+- "url": "https://android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz", -++ "urls": [ -++ "https://mirror.bazel.build/android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz" -++ ], -++ "sha256": "12e44fdbd219c5e1cc62099c2a01d775957603d2d4f693f8285f9d95d9a04e77", -+ "build_file": "@@rules_android+//bzlmod_extensions:apksig.BUILD" -+ } -+ } -+@@ -405,8 +547,8 @@ -+ }, -+ "@@rules_android+//bzlmod_extensions:com_android_dex.bzl%com_android_dex_extension": { -+ "general": { -+- "bzlTransitiveDigest": "rvWbJQc8jInfIAaXIMhSOqUlwM9HVeLey6q0ISvg08Y=", -+- "usagesDigest": "toF8IFMu98H/VU2p1sfVC5fVXVYJunpbbmtM6tOsQXY=", -++ "bzlTransitiveDigest": "K0jbWcRwfM8njdIXNRjRvdApKmBfKeFLScDH+5LSSE0=", -++ "usagesDigest": "0hluQmaWiWak6sVMP5L4wXhNyIwv9fw0y5JJ8lnPb1c=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+@@ -414,8 +556,11 @@ -+ "com_android_dex": { -+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -+ "attributes": { -+- "url": "https://android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz", -+- "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD" -++ "urls": [ -++ "https://mirror.bazel.build/android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz" -++ ], -++ "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD", -++ "sha256": "86b4848c038bf687fadc812239cb01fb8d1d15cef3125b480a0448360992b95d" -+ } -+ } -+ }, -+@@ -444,10 +589,144 @@ -+ "recordedRepoMappingEntries": [] -+ } -+ }, -++ "@@rules_nodejs+//nodejs:extensions.bzl%node": { -++ "general": { -++ "bzlTransitiveDigest": "4pUxCNc22K4I+6+4Nxu52Hur12tFRfa1JMsN5mdDv60=", -++ "usagesDigest": "dqOjZvNvw6/DVBPAiKrXJNA0Tx4GT4Vj/VdUyGMpDL8=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": {}, -++ "generatedRepoSpecs": { -++ "nodejs_linux_amd64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "linux_amd64" -++ } -++ }, -++ "nodejs_linux_arm64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "linux_arm64" -++ } -++ }, -++ "nodejs_linux_s390x": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "linux_s390x" -++ } -++ }, -++ "nodejs_linux_ppc64le": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "linux_ppc64le" -++ } -++ }, -++ "nodejs_darwin_amd64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "darwin_amd64" -++ } -++ }, -++ "nodejs_darwin_arm64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "darwin_arm64" -++ } -++ }, -++ "nodejs_windows_amd64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "windows_amd64" -++ } -++ }, -++ "nodejs_windows_arm64": { -++ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", -++ "attributes": { -++ "node_download_auth": {}, -++ "node_repositories": {}, -++ "node_urls": [ -++ "https://nodejs.org/dist/v{version}/{filename}" -++ ], -++ "node_version": "20.18.0", -++ "include_headers": false, -++ "platform": "windows_arm64" -++ } -++ }, -++ "nodejs": { -++ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", -++ "attributes": { -++ "user_node_repository_name": "nodejs" -++ } -++ }, -++ "nodejs_host": { -++ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", -++ "attributes": { -++ "user_node_repository_name": "nodejs" -++ } -++ }, -++ "nodejs_toolchains": { -++ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_toolchains_repo.bzl%nodejs_toolchains_repo", -++ "attributes": { -++ "user_node_repository_name": "nodejs" -++ } -++ } -++ }, -++ "recordedRepoMappingEntries": [] -++ } -++ }, -+ "@@rules_python+//python/extensions:config.bzl%config": { -+ "general": { -+ "bzlTransitiveDigest": "EcMcbtKZvYmd5Mi1Fpg4EeBBztLHEE5tjO5tLDBYDuU=", -+- "usagesDigest": "EocbSr4I3/Shk4QaFool8b8navUiUFqgzF9bOaaYfFk=", -++ "usagesDigest": "p2al+dDKI5UlCyNvheMVynbWSGbdiji/jMz53fMNfJA=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+@@ -682,7 +961,7 @@ -+ "@@rules_python+//python/uv:uv.bzl%uv": { -+ "general": { -+ "bzlTransitiveDigest": "ijW9KS7qsIY+yBVvJ+Nr1mzwQox09j13DnE3iIwaeTM=", -+- "usagesDigest": "yXvWfXAzpBeW71mWgwU3AqAzXX/dFACnx12eYvBsJ8w=", -++ "usagesDigest": "/HRt5Hw/vpDr9CDrKEPjeDIjxo4307VLxMu8BNAEDWA=", -+ "recordedFileInputs": {}, -+ "recordedDirentsInputs": {}, -+ "envVariables": {}, -+@@ -719,6 +998,533 @@ -+ ] -+ ] -+ } -++ }, -++ "@@rules_rust+//crate_universe:extension.bzl%crate": { -++ "general": { -++ "bzlTransitiveDigest": "VVbU93QvGxFMzb9BcpYTYyyYDpj10Ya6Zm5RH1JEUhw=", -++ "usagesDigest": "EuFUqVKVHF263jHTWOHXs4tFACdRNVOhwpoytdk19bs=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": { -++ "CARGO_BAZEL_DEBUG": null, -++ "CARGO_BAZEL_GENERATOR_SHA256": null, -++ "CARGO_BAZEL_GENERATOR_URL": null, -++ "CARGO_BAZEL_ISOLATED": null, -++ "CARGO_BAZEL_REPIN": null, -++ "CARGO_BAZEL_REPIN_ONLY": null, -++ "CARGO_BAZEL_TIMEOUT": null, -++ "REPIN": null -++ }, -++ "generatedRepoSpecs": { -++ "crates": { -++ "repoRuleId": "@@rules_rust+//crate_universe:extensions.bzl%_generate_repo", -++ "attributes": { -++ "contents": { -++ "BUILD.bazel": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\npackage(default_visibility = [\"//visibility:public\"])\n\nexports_files(\n [\n \"cargo-bazel.json\",\n \"crates.bzl\",\n \"defs.bzl\",\n ] + glob(\n allow_empty = True,\n include = [\"*.bazel\"],\n ),\n)\n\nfilegroup(\n name = \"srcs\",\n srcs = glob(\n allow_empty = True,\n include = [\n \"*.bazel\",\n \"*.bzl\",\n ],\n ),\n)\n\n# Workspace Member Dependencies\nalias(\n name = \"googletest-0.14.3\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"googletest\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme-0.3.37\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste-1.0.15\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote-1.0.47\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn-3.0.3\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n", -++ "alias_rules.bzl": "\"\"\"Alias that transitions its target to `compilation_mode=opt`. Use `transition_alias=\"opt\"` to enable.\"\"\"\n\nload(\"@rules_cc//cc:defs.bzl\", \"CcInfo\")\nload(\"@rules_rust//rust:rust_common.bzl\", \"COMMON_PROVIDERS\")\n\ndef _transition_alias_impl(ctx):\n # `ctx.attr.actual` is a list of 1 item due to the transition\n providers = [ctx.attr.actual[0][provider] for provider in COMMON_PROVIDERS]\n if CcInfo in ctx.attr.actual[0]:\n providers.append(ctx.attr.actual[0][CcInfo])\n return providers\n\ndef _change_compilation_mode(compilation_mode):\n def _change_compilation_mode_impl(_settings, _attr):\n return {\n \"//command_line_option:compilation_mode\": compilation_mode,\n }\n\n return transition(\n implementation = _change_compilation_mode_impl,\n inputs = [],\n outputs = [\n \"//command_line_option:compilation_mode\",\n ],\n )\n\ndef _transition_alias_rule(compilation_mode):\n return rule(\n implementation = _transition_alias_impl,\n provides = COMMON_PROVIDERS,\n attrs = {\n \"actual\": attr.label(\n mandatory = True,\n doc = \"`rust_library()` target to transition to `compilation_mode=opt`.\",\n providers = COMMON_PROVIDERS,\n cfg = _change_compilation_mode(compilation_mode),\n ),\n \"_allowlist_function_transition\": attr.label(\n default = \"@bazel_tools//tools/allowlists/function_transition_allowlist\",\n ),\n },\n doc = \"Transitions a Rust library crate to the `compilation_mode=opt`.\",\n )\n\ntransition_alias_dbg = _transition_alias_rule(\"dbg\")\ntransition_alias_fastbuild = _transition_alias_rule(\"fastbuild\")\ntransition_alias_opt = _transition_alias_rule(\"opt\")\n", -++ "defs.bzl": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\"\"\"\n# `crates_repository` API\n\n- [aliases](#aliases)\n- [crate_deps](#crate_deps)\n- [all_crate_deps](#all_crate_deps)\n- [crate_repositories](#crate_repositories)\n\n\"\"\"\n\nload(\"@bazel_tools//tools/build_defs/repo:git.bzl\", \"new_git_repository\")\nload(\"@bazel_tools//tools/build_defs/repo:http.bzl\", \"http_archive\")\nload(\"@bazel_tools//tools/build_defs/repo:utils.bzl\", \"maybe\")\nload(\"@bazel_skylib//lib:selects.bzl\", \"selects\")\nload(\"@rules_rust//crate_universe/private:local_crate_mirror.bzl\", \"local_crate_mirror\")\n\n###############################################################################\n# MACROS API\n###############################################################################\n\n# An identifier that represent common dependencies (unconditional).\n_COMMON_CONDITION = \"\"\n\ndef _flatten_dependency_maps(all_dependency_maps):\n \"\"\"Flatten a list of dependency maps into one dictionary.\n\n Dependency maps have the following structure:\n\n ```python\n DEPENDENCIES_MAP = {\n # The first key in the map is a Bazel package\n # name of the workspace this file is defined in.\n \"workspace_member_package\": {\n\n # Not all dependencies are supported for all platforms.\n # the condition key is the condition required to be true\n # on the host platform.\n \"condition\": {\n\n # An alias to a crate target. # The label of the crate target the\n # Aliases are only crate names. # package name refers to.\n \"package_name\": \"@full//:label\",\n }\n }\n }\n ```\n\n Args:\n all_dependency_maps (list): A list of dicts as described above\n\n Returns:\n dict: A dictionary as described above\n \"\"\"\n dependencies = {}\n\n for workspace_deps_map in all_dependency_maps:\n for pkg_name, conditional_deps_map in workspace_deps_map.items():\n if pkg_name not in dependencies:\n non_frozen_map = dict()\n for key, values in conditional_deps_map.items():\n non_frozen_map.update({key: dict(values.items())})\n dependencies.setdefault(pkg_name, non_frozen_map)\n continue\n\n for condition, deps_map in conditional_deps_map.items():\n # If the condition has not been recorded, do so and continue\n if condition not in dependencies[pkg_name]:\n dependencies[pkg_name].setdefault(condition, dict(deps_map.items()))\n continue\n\n # Alert on any miss-matched dependencies\n inconsistent_entries = []\n for crate_name, crate_label in deps_map.items():\n existing = dependencies[pkg_name][condition].get(crate_name)\n if existing and existing != crate_label:\n inconsistent_entries.append((crate_name, existing, crate_label))\n dependencies[pkg_name][condition].update({crate_name: crate_label})\n\n return dependencies\n\ndef crate_deps(deps, package_name = None):\n \"\"\"Finds the fully qualified label of the requested crates for the package where this macro is called.\n\n Args:\n deps (list): The desired list of crate targets.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()`.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if not deps:\n return []\n\n if package_name == None:\n package_name = native.package_name()\n\n # Join both sets of dependencies\n dependencies = _flatten_dependency_maps([\n _NORMAL_DEPENDENCIES,\n _NORMAL_DEV_DEPENDENCIES,\n _PROC_MACRO_DEPENDENCIES,\n _PROC_MACRO_DEV_DEPENDENCIES,\n _BUILD_DEPENDENCIES,\n _BUILD_PROC_MACRO_DEPENDENCIES,\n ]).pop(package_name, {})\n\n # Combine all conditional packages so we can easily index over a flat list\n # TODO: Perhaps this should actually return select statements and maintain\n # the conditionals of the dependencies\n flat_deps = {}\n for deps_set in dependencies.values():\n for crate_name, crate_label in deps_set.items():\n flat_deps.update({crate_name: crate_label})\n\n missing_crates = []\n crate_targets = []\n for crate_target in deps:\n if crate_target not in flat_deps:\n missing_crates.append(crate_target)\n else:\n crate_targets.append(flat_deps[crate_target])\n\n if missing_crates:\n fail(\"Could not find crates `{}` among dependencies of `{}`. Available dependencies were `{}`\".format(\n missing_crates,\n package_name,\n dependencies,\n ))\n\n return crate_targets\n\ndef all_crate_deps(\n normal = False, \n normal_dev = False, \n proc_macro = False, \n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Finds the fully qualified label of all requested direct crate dependencies \\\n for the package where this macro is called.\n\n If no parameters are set, all normal dependencies are returned. Setting any one flag will\n otherwise impact the contents of the returned list.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list.\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_dependency_maps = []\n if normal:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n if normal_dev:\n all_dependency_maps.append(_NORMAL_DEV_DEPENDENCIES)\n if proc_macro:\n all_dependency_maps.append(_PROC_MACRO_DEPENDENCIES)\n if proc_macro_dev:\n all_dependency_maps.append(_PROC_MACRO_DEV_DEPENDENCIES)\n if build:\n all_dependency_maps.append(_BUILD_DEPENDENCIES)\n if build_proc_macro:\n all_dependency_maps.append(_BUILD_PROC_MACRO_DEPENDENCIES)\n\n # Default to always using normal dependencies\n if not all_dependency_maps:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n\n dependencies = _flatten_dependency_maps(all_dependency_maps).pop(package_name, None)\n\n if not dependencies:\n if dependencies == None:\n fail(\"Tried to get all_crate_deps for package \" + package_name + \" but that package had no Cargo.toml file\")\n else:\n return []\n\n crate_deps = list(dependencies.pop(_COMMON_CONDITION, {}).values())\n for condition, deps in dependencies.items():\n crate_deps += selects.with_or({\n tuple(_CONDITIONS[condition]): deps.values(),\n \"//conditions:default\": [],\n })\n\n return crate_deps\n\ndef aliases(\n normal = False,\n normal_dev = False,\n proc_macro = False,\n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Produces a map of Crate alias names to their original label\n\n If no dependency kinds are specified, `normal` and `proc_macro` are used by default.\n Setting any one flag will otherwise determine the contents of the returned dict.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list..\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n dict: The aliases of all associated packages\n \"\"\"\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_aliases_maps = []\n if normal:\n all_aliases_maps.append(_NORMAL_ALIASES)\n if normal_dev:\n all_aliases_maps.append(_NORMAL_DEV_ALIASES)\n if proc_macro:\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n if proc_macro_dev:\n all_aliases_maps.append(_PROC_MACRO_DEV_ALIASES)\n if build:\n all_aliases_maps.append(_BUILD_ALIASES)\n if build_proc_macro:\n all_aliases_maps.append(_BUILD_PROC_MACRO_ALIASES)\n\n # Default to always using normal aliases\n if not all_aliases_maps:\n all_aliases_maps.append(_NORMAL_ALIASES)\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n\n aliases = _flatten_dependency_maps(all_aliases_maps).pop(package_name, None)\n\n if not aliases:\n return dict()\n\n common_items = aliases.pop(_COMMON_CONDITION, {}).items()\n\n # If there are only common items in the dictionary, immediately return them\n if not len(aliases.keys()) == 1:\n return dict(common_items)\n\n # Build a single select statement where each conditional has accounted for the\n # common set of aliases.\n crate_aliases = {\"//conditions:default\": dict(common_items)}\n for condition, deps in aliases.items():\n condition_triples = _CONDITIONS[condition]\n for triple in condition_triples:\n if triple in crate_aliases:\n crate_aliases[triple].update(deps)\n else:\n crate_aliases.update({triple: dict(deps.items() + common_items)})\n\n return select(crate_aliases)\n\n###############################################################################\n# WORKSPACE MEMBER DEPS AND ALIASES\n###############################################################################\n\n_NORMAL_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"googletest\": Label(\"@crates//:googletest-0.14.3\"),\n \"linkme\": Label(\"@crates//:linkme-0.3.37\"),\n \"quote\": Label(\"@crates//:quote-1.0.47\"),\n \"syn\": Label(\"@crates//:syn-3.0.3\"),\n },\n },\n}\n\n\n_NORMAL_ALIASES = {\n \"\": {\n _COMMON_CONDITION: {\n },\n },\n}\n\n\n_NORMAL_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_NORMAL_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"paste\": Label(\"@crates//:paste-1.0.15\"),\n },\n },\n}\n\n\n_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_CONDITIONS = {\n \"aarch64-apple-darwin\": [\"@rules_rust//rust/platform:aarch64-apple-darwin\"],\n \"aarch64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\"],\n \"wasm32-unknown-unknown\": [\"@rules_rust//rust/platform:wasm32-unknown-unknown\"],\n \"wasm32-wasip1\": [\"@rules_rust//rust/platform:wasm32-wasip1\"],\n \"x86_64-pc-windows-msvc\": [\"@rules_rust//rust/platform:x86_64-pc-windows-msvc\"],\n \"x86_64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\"],\n \"x86_64-unknown-nixos-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\"],\n}\n\n###############################################################################\n\ndef crate_repositories():\n \"\"\"A macro for defining repositories for all generated crates.\n\n Returns:\n A list of repos visible to the module through the module extension.\n \"\"\"\n maybe(\n http_archive,\n name = \"crates__aho-corasick-1.1.5\",\n sha256 = \"c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/aho-corasick/1.1.5/download\"],\n strip_prefix = \"aho-corasick-1.1.5\",\n build_file = Label(\"@crates//crates:BUILD.aho-corasick-1.1.5.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__autocfg-1.5.1\",\n sha256 = \"f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/autocfg/1.5.1/download\"],\n strip_prefix = \"autocfg-1.5.1\",\n build_file = Label(\"@crates//crates:BUILD.autocfg-1.5.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest-0.14.3\",\n sha256 = \"f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest/0.14.3/download\"],\n strip_prefix = \"googletest-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest_macro-0.14.3\",\n sha256 = \"2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest_macro/0.14.3/download\"],\n strip_prefix = \"googletest_macro-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest_macro-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-0.3.37\",\n sha256 = \"3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme/0.3.37/download\"],\n strip_prefix = \"linkme-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-impl-0.3.37\",\n sha256 = \"77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme-impl/0.3.37/download\"],\n strip_prefix = \"linkme-impl-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-impl-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__memchr-2.8.3\",\n sha256 = \"cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/memchr/2.8.3/download\"],\n strip_prefix = \"memchr-2.8.3\",\n build_file = Label(\"@crates//crates:BUILD.memchr-2.8.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__num-traits-0.2.19\",\n sha256 = \"071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/num-traits/0.2.19/download\"],\n strip_prefix = \"num-traits-0.2.19\",\n build_file = Label(\"@crates//crates:BUILD.num-traits-0.2.19.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__paste-1.0.15\",\n sha256 = \"57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/paste/1.0.15/download\"],\n strip_prefix = \"paste-1.0.15\",\n build_file = Label(\"@crates//crates:BUILD.paste-1.0.15.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__proc-macro2-1.0.107\",\n sha256 = \"985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/proc-macro2/1.0.107/download\"],\n strip_prefix = \"proc-macro2-1.0.107\",\n build_file = Label(\"@crates//crates:BUILD.proc-macro2-1.0.107.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__quote-1.0.47\",\n sha256 = \"1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/quote/1.0.47/download\"],\n strip_prefix = \"quote-1.0.47\",\n build_file = Label(\"@crates//crates:BUILD.quote-1.0.47.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-1.13.1\",\n sha256 = \"f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex/1.13.1/download\"],\n strip_prefix = \"regex-1.13.1\",\n build_file = Label(\"@crates//crates:BUILD.regex-1.13.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-automata-0.4.18\",\n sha256 = \"ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-automata/0.4.18/download\"],\n strip_prefix = \"regex-automata-0.4.18\",\n build_file = Label(\"@crates//crates:BUILD.regex-automata-0.4.18.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-syntax-0.8.11\",\n sha256 = \"d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-syntax/0.8.11/download\"],\n strip_prefix = \"regex-syntax-0.8.11\",\n build_file = Label(\"@crates//crates:BUILD.regex-syntax-0.8.11.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__rustversion-1.0.23\",\n sha256 = \"cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/rustversion/1.0.23/download\"],\n strip_prefix = \"rustversion-1.0.23\",\n build_file = Label(\"@crates//crates:BUILD.rustversion-1.0.23.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-2.0.119\",\n sha256 = \"872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/2.0.119/download\"],\n strip_prefix = \"syn-2.0.119\",\n build_file = Label(\"@crates//crates:BUILD.syn-2.0.119.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-3.0.3\",\n sha256 = \"53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/3.0.3/download\"],\n strip_prefix = \"syn-3.0.3\",\n build_file = Label(\"@crates//crates:BUILD.syn-3.0.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__unicode-ident-1.0.24\",\n sha256 = \"e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/unicode-ident/1.0.24/download\"],\n strip_prefix = \"unicode-ident-1.0.24\",\n build_file = Label(\"@crates//crates:BUILD.unicode-ident-1.0.24.bazel\"),\n )\n\n return [\n struct(repo=\"crates__googletest-0.14.3\", is_dev_dep = False),\n struct(repo=\"crates__linkme-0.3.37\", is_dev_dep = False),\n struct(repo=\"crates__paste-1.0.15\", is_dev_dep = False),\n struct(repo=\"crates__quote-1.0.47\", is_dev_dep = False),\n struct(repo=\"crates__syn-3.0.3\", is_dev_dep = False),\n ]\n" -++ } -++ } -++ }, -++ "crates__aho-corasick-1.1.5": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/aho-corasick/1.1.5/download" -++ ], -++ "strip_prefix": "aho-corasick-1.1.5", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"aho_corasick\",\n deps = [\n \"@crates__memchr-2.8.3//:memchr\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"perf-literal\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=aho-corasick\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.1.5\",\n)\n" -++ } -++ }, -++ "crates__autocfg-1.5.1": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/autocfg/1.5.1/download" -++ ], -++ "strip_prefix": "autocfg-1.5.1", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"autocfg\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2015\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=autocfg\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.5.1\",\n)\n" -++ } -++ }, -++ "crates__googletest-0.14.3": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/googletest/0.14.3/download" -++ ], -++ "strip_prefix": "googletest-0.14.3", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"googletest\",\n deps = [\n \"@crates__num-traits-0.2.19//:num_traits\",\n \"@crates__regex-1.13.1//:regex\",\n ],\n proc_macro_deps = [\n \"@crates__googletest_macro-0.14.3//:googletest_macro\",\n \"@crates__rustversion-1.0.23//:rustversion\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" -++ } -++ }, -++ "crates__googletest_macro-0.14.3": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/googletest_macro/0.14.3/download" -++ ], -++ "strip_prefix": "googletest_macro-0.14.3", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"googletest_macro\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-2.0.119//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest_macro\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" -++ } -++ }, -++ "crates__linkme-0.3.37": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/linkme/0.3.37/download" -++ ], -++ "strip_prefix": "linkme-0.3.37", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"linkme\",\n deps = [\n \"@crates__linkme-0.3.37//:build_script_build\",\n ],\n proc_macro_deps = [\n \"@crates__linkme-impl-0.3.37//:linkme_impl\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__linkme-impl-0.3.37": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/linkme-impl/0.3.37/download" -++ ], -++ "strip_prefix": "linkme-impl-0.3.37", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"linkme_impl\",\n deps = [\n \"@crates__linkme-impl-0.3.37//:build_script_build\",\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-3.0.3//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme-impl\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__memchr-2.8.3": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/memchr/2.8.3/download" -++ ], -++ "strip_prefix": "memchr-2.8.3", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"memchr\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=memchr\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.8.3\",\n)\n" -++ } -++ }, -++ "crates__num-traits-0.2.19": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/num-traits/0.2.19/download" -++ ], -++ "strip_prefix": "num-traits-0.2.19", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"num_traits\",\n deps = [\n \"@crates__num-traits-0.2.19//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.2.19\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n deps = [\n \"@crates__autocfg-1.5.1//:autocfg\",\n ],\n edition = \"2021\",\n pkg_name = \"num-traits\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.2.19\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__paste-1.0.15": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/paste/1.0.15/download" -++ ], -++ "strip_prefix": "paste-1.0.15", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"paste\",\n deps = [\n \"@crates__paste-1.0.15//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.15\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"paste\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.15\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__proc-macro2-1.0.107": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/proc-macro2/1.0.107/download" -++ ], -++ "strip_prefix": "proc-macro2-1.0.107", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"proc_macro2\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:build_script_build\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.107\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"proc-macro2\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.107\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__quote-1.0.47": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/quote/1.0.47/download" -++ ], -++ "strip_prefix": "quote-1.0.47", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"quote\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.47\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"quote\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.47\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__regex-1.13.1": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/regex/1.13.1/download" -++ ], -++ "strip_prefix": "regex-1.13.1", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-automata-0.4.18//:regex_automata\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"perf\",\n \"perf-backtrack\",\n \"perf-cache\",\n \"perf-dfa\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-onepass\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.13.1\",\n)\n" -++ } -++ }, -++ "crates__regex-automata-0.4.18": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/regex-automata/0.4.18/download" -++ ], -++ "strip_prefix": "regex-automata-0.4.18", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_automata\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"dfa-onepass\",\n \"hybrid\",\n \"meta\",\n \"nfa-backtrack\",\n \"nfa-pikevm\",\n \"nfa-thompson\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-literal-multisubstring\",\n \"perf-literal-substring\",\n \"std\",\n \"syntax\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n \"unicode-word-boundary\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-automata\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.4.18\",\n)\n" -++ } -++ }, -++ "crates__regex-syntax-0.8.11": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/regex-syntax/0.8.11/download" -++ ], -++ "strip_prefix": "regex-syntax-0.8.11", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_syntax\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-syntax\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.8.11\",\n)\n" -++ } -++ }, -++ "crates__rustversion-1.0.23": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/rustversion/1.0.23/download" -++ ], -++ "strip_prefix": "rustversion-1.0.23", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"rustversion\",\n deps = [\n \"@crates__rustversion-1.0.23//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.23\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build/build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"rustversion\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.23\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" -++ } -++ }, -++ "crates__syn-2.0.119": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/syn/2.0.119/download" -++ ], -++ "strip_prefix": "syn-2.0.119", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"extra-traits\",\n \"full\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.0.119\",\n)\n" -++ } -++ }, -++ "crates__syn-3.0.3": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/syn/3.0.3/download" -++ ], -++ "strip_prefix": "syn-3.0.3", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"3.0.3\",\n)\n" -++ } -++ }, -++ "crates__unicode-ident-1.0.24": { -++ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", -++ "attributes": { -++ "patch_args": [], -++ "patch_tool": "", -++ "patches": [], -++ "remote_patch_strip": 1, -++ "sha256": "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75", -++ "type": "tar.gz", -++ "urls": [ -++ "https://static.crates.io/crates/unicode-ident/1.0.24/download" -++ ], -++ "strip_prefix": "unicode-ident-1.0.24", -++ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"unicode_ident\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=unicode-ident\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.24\",\n)\n" -++ } -++ } -++ }, -++ "recordedRepoMappingEntries": [ -++ [ -++ "bazel_features+", -++ "bazel_features_globals", -++ "bazel_features++version_extension+bazel_features_globals" -++ ], -++ [ -++ "bazel_features+", -++ "bazel_features_version", -++ "bazel_features++version_extension+bazel_features_version" -++ ], -++ [ -++ "rules_cc+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "rules_cc+", -++ "cc_compatibility_proxy", -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -++ ], -++ [ -++ "rules_cc+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_features", -++ "bazel_features+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_skylib", -++ "bazel_skylib+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "rules_rust+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_rust+", -++ "rules_rust", -++ "rules_rust+" -++ ] -++ ] -++ } -++ }, -++ "@@rules_rust+//crate_universe/private:internal_extensions.bzl%cu_nr": { -++ "general": { -++ "bzlTransitiveDigest": "GOOgbXFJQhO4daGipwnspaixIHp6AWTvXRBe2wMULd4=", -++ "usagesDigest": "tG3p3Nb5XxC7vWY/bcKdb//g0HoAxpxxH3F5/jBVlk4=", -++ "recordedFileInputs": {}, -++ "recordedDirentsInputs": {}, -++ "envVariables": {}, -++ "generatedRepoSpecs": { -++ "cargo_bazel_bootstrap": { -++ "repoRuleId": "@@rules_rust+//cargo/private:cargo_bootstrap.bzl%cargo_bootstrap_repository", -++ "attributes": { -++ "srcs": [ -++ "@@rules_rust+//crate_universe:src/api.rs", -++ "@@rules_rust+//crate_universe:src/api/lockfile.rs", -++ "@@rules_rust+//crate_universe:src/cli.rs", -++ "@@rules_rust+//crate_universe:src/cli/generate.rs", -++ "@@rules_rust+//crate_universe:src/cli/query.rs", -++ "@@rules_rust+//crate_universe:src/cli/render.rs", -++ "@@rules_rust+//crate_universe:src/cli/splice.rs", -++ "@@rules_rust+//crate_universe:src/cli/vendor.rs", -++ "@@rules_rust+//crate_universe:src/config.rs", -++ "@@rules_rust+//crate_universe:src/context.rs", -++ "@@rules_rust+//crate_universe:src/context/crate_context.rs", -++ "@@rules_rust+//crate_universe:src/context/platforms.rs", -++ "@@rules_rust+//crate_universe:src/lib.rs", -++ "@@rules_rust+//crate_universe:src/lockfile.rs", -++ "@@rules_rust+//crate_universe:src/main.rs", -++ "@@rules_rust+//crate_universe:src/metadata.rs", -++ "@@rules_rust+//crate_universe:src/metadata/cargo_bin.rs", -++ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_resolver.rs", -++ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.bat", -++ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.sh", -++ "@@rules_rust+//crate_universe:src/metadata/dependency.rs", -++ "@@rules_rust+//crate_universe:src/metadata/metadata_annotation.rs", -++ "@@rules_rust+//crate_universe:src/rendering.rs", -++ "@@rules_rust+//crate_universe:src/rendering/template_engine.rs", -++ "@@rules_rust+//crate_universe:src/rendering/templates/module_bzl.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/partials/header.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/aliases_map.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/deps_map.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_git.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_http.j2", -++ "@@rules_rust+//crate_universe:src/rendering/templates/vendor_module.j2", -++ "@@rules_rust+//crate_universe:src/rendering/verbatim/alias_rules.bzl", -++ "@@rules_rust+//crate_universe:src/select.rs", -++ "@@rules_rust+//crate_universe:src/splicing.rs", -++ "@@rules_rust+//crate_universe:src/splicing/cargo_config.rs", -++ "@@rules_rust+//crate_universe:src/splicing/crate_index_lookup.rs", -++ "@@rules_rust+//crate_universe:src/splicing/splicer.rs", -++ "@@rules_rust+//crate_universe:src/test.rs", -++ "@@rules_rust+//crate_universe:src/utils.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/glob.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/label.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/select.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/select_dict.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/select_list.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/select_scalar.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/select_set.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/serialize.rs", -++ "@@rules_rust+//crate_universe:src/utils/starlark/target_compatible_with.rs", -++ "@@rules_rust+//crate_universe:src/utils/symlink.rs", -++ "@@rules_rust+//crate_universe:src/utils/target_triple.rs" -++ ], -++ "binary": "cargo-bazel", -++ "cargo_lockfile": "@@rules_rust+//crate_universe:Cargo.lock", -++ "cargo_toml": "@@rules_rust+//crate_universe:Cargo.toml", -++ "version": "1.93.1", -++ "timeout": 900, -++ "rust_toolchain_cargo_template": "@rust_host_tools//:bin/{tool}", -++ "rust_toolchain_rustc_template": "@rust_host_tools//:bin/{tool}", -++ "compressed_windows_toolchain_names": false -++ } -++ } -++ }, -++ "moduleExtensionMetadata": { -++ "explicitRootModuleDirectDeps": [ -++ "cargo_bazel_bootstrap" -++ ], -++ "explicitRootModuleDirectDevDeps": [], -++ "useAllRepos": "NO", -++ "reproducible": false -++ }, -++ "recordedRepoMappingEntries": [ -++ [ -++ "bazel_features+", -++ "bazel_features_globals", -++ "bazel_features++version_extension+bazel_features_globals" -++ ], -++ [ -++ "bazel_features+", -++ "bazel_features_version", -++ "bazel_features++version_extension+bazel_features_version" -++ ], -++ [ -++ "rules_cc+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "rules_cc+", -++ "cc_compatibility_proxy", -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy" -++ ], -++ [ -++ "rules_cc+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_cc++compatibility_proxy+cc_compatibility_proxy", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_features", -++ "bazel_features+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_skylib", -++ "bazel_skylib+" -++ ], -++ [ -++ "rules_rust+", -++ "bazel_tools", -++ "bazel_tools" -++ ], -++ [ -++ "rules_rust+", -++ "cui", -++ "rules_rust++cu+cui" -++ ], -++ [ -++ "rules_rust+", -++ "rrc", -++ "rules_rust++i2+rrc" -++ ], -++ [ -++ "rules_rust+", -++ "rules_cc", -++ "rules_cc+" -++ ], -++ [ -++ "rules_rust+", -++ "rules_rust", -++ "rules_rust+" -++ ] -++ ] -++ } -+ } -+ }, -+ "facts": { +-+ // Unlock the original locked mutex. +-+ // Calling try_lock() resolves the underlying mutex and re-enters it on this +-+ // thread (incrementing the recursive count), so we unlock twice to fully +-+ // release both acquisitions. +-+ TargetAPIMutex cleanup_lock(target_sp); +-+ ASSERT_TRUE(cleanup_lock.try_lock()); +-+ cleanup_lock.unlock(); +-+ cleanup_lock.unlock(); +- } +- +- TEST_F(TargetAPIMutexTargetTest, LockGuardReleasesOnScopeExit) { + diff -ruN --strip-trailing-cr a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp + --- a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp + +++ b/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp +-@@ -5937,6 +5937,8 @@ ++@@ -5964,6 +5964,8 @@ + ISD::matchUnaryPredicate( + Y, + [&](auto *C) { +@@ -133,72 +41,6 @@ diff -ruN --strip-trailing-cr a/llvm/lib/CodeGen/SelectionDAG/DAGCombiner.cpp b/ + const APInt &YConst = C->getAsAPIntVal(); + return (Opcode == ISD::ABDS) + ? YConst.isSignedIntN(Bits) +-diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp +---- a/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp +-+++ b/llvm/lib/Target/AArch64/AArch64ISelDAGToDAG.cpp +-@@ -1254,18 +1254,6 @@ +- return true; +- } +- +--/// Check whether \p GAN is the low part of a TLS address computation, i.e. the +--/// second operand of an ADDlow. The target flags on their own do not tell the +--/// ELF local-exec (:tprel_lo12: and :tprel_lo12_nc:) cases apart from other +--/// uses, so callers that depend on local-exec semantics have to check the +--/// object format as well. Local dynamic never gets here because it does not +--/// build an ADDlow. +--static bool isTLSLo12(const GlobalAddressSDNode *GAN) { +-- unsigned Flags = GAN->getTargetFlags(); +-- return (Flags & (AArch64II::MO_TLS | AArch64II::MO_FRAGMENT)) == +-- (AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); +--} +-- +- /// Check if the immediate offset is valid as a scaled immediate. +- static bool isValidAsScaledImmediate(int64_t Offset, unsigned Range, +- unsigned Size) { +-@@ -1361,13 +1349,8 @@ +- if (!GAN) +- return true; +- +-- // Folding the low part of an ELF local-exec TLS address into a 128-bit +-- // access needs R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC variant, which +-- // the GNU bfd linker does not support, so keep materialising the address +-- // with an add. +- if (GAN->getOffset() % Size == 0 && +-- GAN->getGlobal()->getPointerAlignment(DL) >= Size && +-- !(Size > 8 && Subtarget->isTargetELF() && isTLSLo12(GAN))) +-+ GAN->getGlobal()->getPointerAlignment(DL) >= Size) +- return true; +- } +- +-diff -ruN --strip-trailing-cr a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp +---- a/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp +-+++ b/llvm/lib/Target/AArch64/AArch64ISelLowering.cpp +-@@ -11542,7 +11542,10 @@ +- // add x0, x0, :tprel_lo12:a +- SDValue Var = DAG.getTargetGlobalAddress( +- GV, DL, PtrVT, 0, AArch64II::MO_TLS | AArch64II::MO_PAGEOFF); +-- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, ThreadBase, Var); +-+ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, ThreadBase, +-+ Var, +-+ DAG.getTargetConstant(0, DL, MVT::i32)), +-+ 0); +- } +- +- case 24: { +-@@ -11558,10 +11561,9 @@ +- HiVar, +- DAG.getTargetConstant(0, DL, MVT::i32)), +- 0); +-- // Emit the low part as an ADDlow so that it can be folded into the +-- // addressing mode of a following load or store, turning the add into a +-- // :tprel_lo12_nc: relocation on the memory access itself. +-- return DAG.getNode(AArch64ISD::ADDlow, DL, PtrVT, Addr, LoVar); +-+ return SDValue(DAG.getMachineNode(AArch64::ADDXri, DL, PtrVT, Addr, LoVar, +-+ DAG.getTargetConstant(0, DL, MVT::i32)), +-+ 0); +- } +- +- case 32: { + diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/abd-combine.ll b/llvm/test/CodeGen/AArch64/abd-combine.ll + --- a/llvm/test/CodeGen/AArch64/abd-combine.ll + +++ b/llvm/test/CodeGen/AArch64/abd-combine.ll +@@ -229,1276 +71,3 @@ diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/abd-combine.ll b/llvm/ + declare <8 x i16> @llvm.aarch64.neon.sabd.v8i16(<8 x i16>, <8 x i16>) + declare <8 x i32> @llvm.abs.v8i32(<8 x i32>, i1) + +declare <4 x i32> @llvm.abs.v4i32(<4 x i32>, i1) +-diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll +---- a/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll +-+++ b/llvm/test/CodeGen/AArch64/arm64-tls-local-exec.ll +-@@ -27,24 +27,24 @@ +- ; RUN: llc -mtriple=arm64-none-linux-gnu -filetype=obj < %s -code-model=large | llvm-objdump -r - | FileCheck --check-prefix=CHECK-24-RELOC %s +- +- @local_exec_var = thread_local(localexec) global i32 0 +--@local_exec_var64 = thread_local(localexec) global i64 0 +--@vec_local_exec_var = thread_local(localexec) global <2 x i64> zeroinitializer, align 16 +- +- define i32 @test_local_exec() { +- ; CHECK-LABEL: test_local_exec: +- %val = load i32, ptr @local_exec_var +- +- ; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 +--; CHECK-12: ldr w0, [x[[R1]], :tprel_lo12:local_exec_var] +-+; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:local_exec_var +-+; CHECK-12: ldr w0, [x[[R2]]] +- +--; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12 +-+; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 +- +- ; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 +- ; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var +--; CHECK-24: ldr w0, [x[[R2]], :tprel_lo12_nc:local_exec_var] +-+; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:local_exec_var +-+; CHECK-24: ldr w0, [x[[R3]]] +- +- ; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 +--; CHECK-24-RELOC: R_AARCH64_TLSLE_LDST32_TPREL_LO12_NC +-+; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC +- +- ; CHECK-32: movz x[[R2:[0-9]+]], #:tprel_g1:local_exec_var +- ; CHECK-32: mrs x[[R1:[0-9]+]], TPIDR_EL0 +-@@ -66,24 +66,6 @@ +- ret i32 %val +- } +- +--define void @test_local_exec_store64(i64 %val) { +--; CHECK-LABEL: test_local_exec_store64: +-- store i64 %val, ptr @local_exec_var64 +-- +--; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 +--; CHECK-12: str x0, [x[[R1]], :tprel_lo12:local_exec_var64] +-- +--; CHECK-12-RELOC: R_AARCH64_TLSLE_LDST64_TPREL_LO12 +-- +--; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 +--; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:local_exec_var64 +--; CHECK-24: str x0, [x[[R2]], :tprel_lo12_nc:local_exec_var64] +-- +--; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 local_exec_var64 +--; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_LDST64_TPREL_LO12_NC local_exec_var64 +-- ret void +--} +-- +- define ptr @test_local_exec_addr() { +- ; CHECK-LABEL: test_local_exec_addr: +- ret ptr @local_exec_var +-@@ -122,26 +104,3 @@ +- ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G1_NC +- ; CHECK-48-RELOC: R_AARCH64_TLSLE_MOVW_TPREL_G0_NC +- } +-- +--; A 128-bit access would need R_AARCH64_TLSLE_LDST128_TPREL_LO12 or its NC +--; variant, which not every linker implements, so the low part stays in a +--; separate add. +--define <2 x i64> @test_local_exec_128bit() { +--; CHECK-LABEL: test_local_exec_128bit: +-- %val = load <2 x i64>, ptr @vec_local_exec_var +-- +--; CHECK-12: mrs x[[R1:[0-9]+]], TPIDR_EL0 +--; CHECK-12: add x[[R2:[0-9]+]], x[[R1]], :tprel_lo12:vec_local_exec_var +--; CHECK-12: ldr q0, [x[[R2]]] +-- +--; CHECK-12-RELOC: R_AARCH64_TLSLE_ADD_TPREL_LO12 vec_local_exec_var +-- +--; CHECK-24: mrs x[[R1:[0-9]+]], TPIDR_EL0 +--; CHECK-24: add x[[R2:[0-9]+]], x[[R1]], :tprel_hi12:vec_local_exec_var +--; CHECK-24: add x[[R3:[0-9]+]], x[[R2]], :tprel_lo12_nc:vec_local_exec_var +--; CHECK-24: ldr q0, [x[[R3]]] +-- +--; CHECK-24-RELOC: R_AARCH64_TLSLE_ADD_TPREL_HI12 vec_local_exec_var +--; CHECK-24-RELOC-NEXT: R_AARCH64_TLSLE_ADD_TPREL_LO12_NC vec_local_exec_var +-- ret <2 x i64> %val +--} +-diff -ruN --strip-trailing-cr a/llvm/test/CodeGen/AArch64/win-tls.ll b/llvm/test/CodeGen/AArch64/win-tls.ll +---- a/llvm/test/CodeGen/AArch64/win-tls.ll +-+++ b/llvm/test/CodeGen/AArch64/win-tls.ll +-@@ -3,7 +3,6 @@ +- @tlsVar = thread_local global i32 0 +- @tlsVar8 = thread_local global i8 0 +- @tlsVar64 = thread_local global i64 0 +--@tlsVar128 = thread_local global <2 x i64> zeroinitializer +- +- define i32 @getVar() { +- %1 = load i32, ptr @tlsVar +-@@ -29,11 +28,6 @@ +- ret i64 %1 +- } +- +--define <2 x i64> @getVar128() { +-- %1 = load <2 x i64>, ptr @tlsVar128 +-- ret <2 x i64> %1 +--} +-- +- ; CHECK-LABEL: getVar +- ; CHECK: adrp [[TLS_INDEX_ADDR:x[0-9]+]], _tls_index +- ; CHECK: ldr [[TLS_POINTER:x[0-9]+]], [x18, #88] +-@@ -68,7 +62,3 @@ +- ; CHECK-LABEL: getVar64 +- ; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar64 +- ; CHECK: ldr x0, [[[TLS]], :secrel_lo12:tlsVar64] +-- +--; CHECK-LABEL: getVar128 +--; CHECK: add [[TLS:x[0-9]+]], [[TLS]], :secrel_hi12:tlsVar128 +--; CHECK: ldr q0, [[[TLS]], :secrel_lo12:tlsVar128] +-diff -ruN --strip-trailing-cr a/utils/bazel/.bazelrc b/utils/bazel/.bazelrc +---- a/utils/bazel/.bazelrc +-+++ b/utils/bazel/.bazelrc +-@@ -222,6 +222,19 @@ +- build:hermetic-toolchain --copt=-Wno-modules-import-nested-redundant --host_copt=-Wno-modules-import-nested-redundant +- +- ############################################################################### +-+# Options for Emscripten WebAssembly builds. +-+############################################################################### +-+ +-+build:wasm --platforms=@emsdk//:platform_wasm +-+# Match LLVM's single-threaded config and shut the runtime down when main returns. +-+build:wasm --features=-use_pthreads,exit_runtime +-+# Make the default CLI artifact use Node's host filesystem; browser builds override this to 0. +-+build:wasm --linkopt=-sNODERAWFS=1 +-+ +-+# TODO: zstd unconditionally enables pthreads on non-Windows targets. +-+build:wasm --@llvm-project//third-party:llvm_enable_zstd=false +-+ +-+############################################################################### +- # Options for continuous integration. +- ############################################################################### +- +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel +---- a/utils/bazel/llvm-project-overlay/clang/BUILD.bazel +-+++ b/utils/bazel/llvm-project-overlay/clang/BUILD.bazel +-@@ -1957,6 +1957,7 @@ +- ":parse", +- ":sema", +- ":serialization", +-+ "//compiler-rt:emutls", +- "//llvm:AllTargetsAsmParsers", +- "//llvm:AllTargetsCodeGens", +- "//llvm:Core", +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel +---- a/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel +-+++ b/utils/bazel/llvm-project-overlay/compiler-rt/BUILD.bazel +-@@ -165,6 +165,15 @@ +- ], +- ) +- +-+cc_library( +-+ name = "emutls", +-+ srcs = [ +-+ "lib/builtins/emutls.c", +-+ ], +-+ hdrs = glob(["lib/builtins/*.h"]), +-+ linkstatic = True, +-+) +-+ +- filegroup( +- name = "fuzzer_installed_hdrs", +- srcs = glob(["include/fuzzer/*.h"]), +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +---- a/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +-+++ b/utils/bazel/llvm-project-overlay/libc/BUILD.bazel +-@@ -1467,15 +1467,12 @@ +- name = "__support_libc_assert", +- hdrs = ["src/__support/libc_assert.h"], +- deps = [ +-- ":__support_integer_to_string", +- ":__support_macros_attributes", +- ":__support_macros_config", +- ":__support_macros_hardening", +- ":__support_macros_macro_utils", +- ":__support_macros_optimization", +- ":__support_macros_properties_os", +-- ":__support_osutil_exit_hdrs", +-- ":__support_osutil_io", +- ], +- ) +- +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl +---- a/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl +-+++ b/utils/bazel/llvm-project-overlay/libc/test/libc_test_rules.bzl +-@@ -66,7 +66,7 @@ +- copts = copts + _FULL_BUILD_COPTS +- +- # Temporarily disable full_build tests (currently broken) to unblock CI. +-- tags = tags + ["manual", "notap"] +-+ tags = tags + ["manual", "nobuildkite", "notap"] +- cc_test( +- name = name, +- local_defines = local_defines + _TEST_DEFINES + LIBC_CONFIGURE_OPTIONS, +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel +---- a/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel +-+++ b/utils/bazel/llvm-project-overlay/llvm/BUILD.bazel +-@@ -456,6 +456,7 @@ +- "-pthread", +- "-ldl", +- ], +-+ "@platforms//os:emscripten": [], +- "//conditions:default": [ +- "-pthread", +- "-ldl", +-@@ -1439,7 +1440,10 @@ +- "include/llvm/Analysis/Utils/*.h", +- ], +- ) + ["include/llvm-c/Analysis.h"], +-- copts = llvm_copts + ["-ftrapping-math"], +-+ copts = llvm_copts + select({ +-+ "@platforms//os:emscripten": [], +-+ "//conditions:default": ["-ftrapping-math"], +-+ }), +- features = ["-parse_headers"], +- textual_hdrs = glob([ +- "include/llvm/Analysis/*.def", +-@@ -4973,6 +4977,7 @@ +- # ll scripts rely on symbols from dependent +- # libraries being resolvable. +- linkopts = select({ +-+ "@platforms//os:emscripten": [], +- "@platforms//os:macos": [], +- "@platforms//os:windows": [], +- "//conditions:default": [ +-@@ -5548,6 +5553,7 @@ +- copts = llvm_copts, +- # Make symbols from the standard library dynamically resolvable. +- linkopts = select({ +-+ "@platforms//os:emscripten": [], +- "@platforms//os:macos": [], +- "@platforms//os:windows": [], +- "//conditions:default": [ +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/config.bzl b/utils/bazel/llvm-project-overlay/llvm/config.bzl +---- a/utils/bazel/llvm-project-overlay/llvm/config.bzl +-+++ b/utils/bazel/llvm-project-overlay/llvm/config.bzl +-@@ -46,7 +46,28 @@ +- "HAVE_UNISTD_H=1", +- ] +- +-+emscripten_defines = [ +-+ "LLVM_ON_UNIX=1", +-+ r'LTDL_SHLIB_EXT=\".so\"', +-+ r'LLVM_PLUGIN_EXT=\".so\"', +-+ "LLVM_ENABLE_LLVM_EXPORT_ANNOTATIONS=1", +-+ "LLVM_ENABLE_PLUGINS=0", +-+ "LLVM_ENABLE_THREADS=0", +-+ "HAVE_MALLINFO=1", +-+ "HAVE_SETENV_R=1", +-+ "HAVE_STRERROR_R=1", +-+ "HAVE_SYSEXITS_H=1", +-+ "HAVE_SYS_IOCTL_H=1", +-+ "HAVE_UNISTD_H=1", +-+] +-+ +-+fenv_defines = [ +-+ "HAVE_DECL_FE_ALL_EXCEPT=1", +-+ "HAVE_DECL_FE_INEXACT=1", +-+] +-+ +- backtrace_defines = select({ +-+ "@platforms//os:emscripten": [], +- "@platforms//os:windows": [], +- "@llvm//platforms/config:musl": [], +- "//conditions:default": [ +-@@ -60,14 +81,14 @@ +- "//conditions:default": [], +- }) +- +--linux_defines = posix_defines + [ +-+linux_defines = posix_defines + fenv_defines + [ +- "_GNU_SOURCE", +- "HAVE_GETAUXVAL=1", +- "HAVE_SBRK=1", +- "HAVE_STRUCT_STAT_ST_MTIM_TV_NSEC=1", +- ] +- +--macos_defines = posix_defines + [ +-+macos_defines = posix_defines + fenv_defines + [ +- "HAVE_MACH_MACH_H=1", +- "HAVE_MALLOC_MALLOC_H=1", +- "HAVE_MALLOC_ZONE_STATISTICS=1", +-@@ -89,12 +110,13 @@ +- # LLVM features +- r'LTDL_SHLIB_EXT=\".dll\"', +- r'LLVM_PLUGIN_EXT=\".dll\"', +--] +-+] + fenv_defines +- +- # TODO: We should switch to platforms-based config settings to make this easier +- # to express. +- os_defines = select({ +-- "@platforms//os:freebsd": posix_defines, +-+ "@platforms//os:emscripten": emscripten_defines, +-+ "@platforms//os:freebsd": posix_defines + fenv_defines, +- "@platforms//os:macos": macos_defines, +- "@platforms//os:windows": win32_defines, +- "//conditions:default": linux_defines, +-@@ -117,6 +139,7 @@ +- Label("//llvm:linux_ppc64le"): native_arch_defines("PowerPC", "powerpc64le-unknown-linux-gnu"), +- Label("//llvm:linux_riscv64"): native_arch_defines("RISCV", "riscv64-unknown-linux-gnu"), +- Label("//llvm:linux_s390x"): native_arch_defines("SystemZ", "systemz-unknown-linux_gnu"), +-+ "@platforms//os:emscripten": native_arch_defines("WebAssembly", "wasm32-unknown-emscripten"), +- "@platforms//os:windows": native_arch_defines("X86", "x86_64-pc-win32"), +- "//conditions:default": native_arch_defines("X86", "x86_64-unknown-linux-gnu"), +- }) + [ +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h +---- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h +-+++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/config.h +-@@ -56,11 +56,11 @@ +- +- /* Define to 1 if you have the declaration of `FE_ALL_EXCEPT', and to 0 if you +- don't. */ +--#define HAVE_DECL_FE_ALL_EXCEPT 1 +-+/* HAVE_DECL_FE_ALL_EXCEPT defined in Bazel */ +- +- /* Define to 1 if you have the declaration of `FE_INEXACT', and to 0 if you +- don't. */ +--#define HAVE_DECL_FE_INEXACT 1 +-+/* HAVE_DECL_FE_INEXACT defined in Bazel */ +- +- /* Define to 1 if you have the declaration of `strerror_s', and to 0 if you +- don't. */ +-diff -ruN --strip-trailing-cr a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h +---- a/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h +-+++ b/utils/bazel/llvm-project-overlay/llvm/include/llvm/Config/llvm-config.h +-@@ -28,7 +28,7 @@ +- /* LLVM_DEFAULT_TARGET_TRIPLE defined in Bazel */ +- +- /* Define if threads enabled */ +--#define LLVM_ENABLE_THREADS 1 +-+/* LLVM_ENABLE_THREADS defined in Bazel */ +- +- /* Has gcc/MSVC atomic intrinsics */ +- #define LLVM_HAS_ATOMICS 1 +-diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel b/utils/bazel/MODULE.bazel +---- a/utils/bazel/MODULE.bazel +-+++ b/utils/bazel/MODULE.bazel +-@@ -30,6 +30,7 @@ +- bazel_dep(name = "libpfm", version = "4.13.0", repo_name = "pfm") +- bazel_dep(name = "vulkan_headers", version = "1.4.349") +- +-+bazel_dep(name = "emsdk", version = "6.0.2", dev_dependency = True) +- bazel_dep(name = "llvm", version = "0.8.5", dev_dependency = True) +- +- llvm_repos_extension = use_extension(":extensions.bzl", "llvm_repos_extension") +-diff -ruN --strip-trailing-cr a/utils/bazel/MODULE.bazel.lock b/utils/bazel/MODULE.bazel.lock +---- a/utils/bazel/MODULE.bazel.lock +-+++ b/utils/bazel/MODULE.bazel.lock +-@@ -81,6 +81,8 @@ +- "https://bcr.bazel.build/modules/buildozer/7.1.2/source.json": "c9028a501d2db85793a6996205c8de120944f50a0d570438fcae0457a5f9d1f8", +- "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/MODULE.bazel": "f6561baff0fc0035c9c1a9e2b0820de106cdb01b37bf5c81276860ccc863e5b2", +- "https://bcr.bazel.build/modules/eigen/3.4.0.bcr.3/source.json": "a8611a2b5577929ad7e1f44ded19dab21a188125a74ac6192d21d283609f280f", +-+ "https://bcr.bazel.build/modules/emsdk/6.0.2/MODULE.bazel": "4a3c4195e5f2e0056bc18bf9f8af631c4f720c3e8cea45cfb7247eacf02e27fe", +-+ "https://bcr.bazel.build/modules/emsdk/6.0.2/source.json": "53111cbcb9f0971aa14da976bafbb937a1bc92a2580c8881890983cd2d289af0", +- "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/MODULE.bazel": "f1b7bb2dd53e8f2ef984b39485ec8a44e9076dda5c4b8efd2fb4c6a6e856a31d", +- "https://bcr.bazel.build/modules/gawk/5.3.2.bcr.3/source.json": "ebe931bfe362e4b41e59ee00a528db6074157ff2ced92eb9e970acab2e1089c9", +- "https://bcr.bazel.build/modules/gazelle/0.32.0/MODULE.bazel": "b499f58a5d0d3537f3cf5b76d8ada18242f64ec474d8391247438bf04f58c7b8", +-@@ -184,6 +186,7 @@ +- "https://bcr.bazel.build/modules/rules_cc/0.2.13/MODULE.bazel": "eecdd666eda6be16a8d9dc15e44b5c75133405e820f620a234acc4b1fdc5aa37", +- "https://bcr.bazel.build/modules/rules_cc/0.2.14/MODULE.bazel": "353c99ed148887ee89c54a17d4100ae7e7e436593d104b668476019023b58df8", +- "https://bcr.bazel.build/modules/rules_cc/0.2.15/MODULE.bazel": "6a0a4a75a57aa6dc888300d848053a58c6b12a29f89d4304e1c41448514ec6e8", +-+ "https://bcr.bazel.build/modules/rules_cc/0.2.16/MODULE.bazel": "9242fa89f950c6ef7702801ab53922e99c69b02310c39fb6e62b2bd30df2a1d4", +- "https://bcr.bazel.build/modules/rules_cc/0.2.17/MODULE.bazel": "1849602c86cb60da8613d2de887f9566a6d354a6df6d7009f9d04a14402f9a84", +- "https://bcr.bazel.build/modules/rules_cc/0.2.18/MODULE.bazel": "4460ec36adc8f722a6a2a4ac9374cb91f2acebadaa93fc37966129afb3dece87", +- "https://bcr.bazel.build/modules/rules_cc/0.2.19/MODULE.bazel": "d5e0f05b63273281a16654eb6b1a8742a75ec153ac8b4f0419949d6e401e46f0", +-@@ -242,6 +245,8 @@ +- "https://bcr.bazel.build/modules/rules_license/0.0.7/MODULE.bazel": "088fbeb0b6a419005b89cf93fe62d9517c0a2b8bb56af3244af65ecfe37e7d5d", +- "https://bcr.bazel.build/modules/rules_license/1.0.0/MODULE.bazel": "a7fda60eefdf3d8c827262ba499957e4df06f659330bbe6cdbdb975b768bb65c", +- "https://bcr.bazel.build/modules/rules_license/1.0.0/source.json": "a52c89e54cc311196e478f8382df91c15f7a2bfdf4c6cd0e2675cc2ff0b56efb", +-+ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/MODULE.bazel": "c22a48b2a0dbf05a9dc5f83837bbc24c226c1f6e618de3c3a610044c9f336056", +-+ "https://bcr.bazel.build/modules/rules_nodejs/6.7.3/source.json": "a3f966f4415a8a6545e560ee5449eac95cc633f96429d08e87c87775c72f5e09", +- "https://bcr.bazel.build/modules/rules_pkg/0.7.0/MODULE.bazel": "df99f03fc7934a4737122518bb87e667e62d780b610910f0447665a7e2be62dc", +- "https://bcr.bazel.build/modules/rules_pkg/1.0.1/MODULE.bazel": "5b1df97dbc29623bccdf2b0dcd0f5cb08e2f2c9050aab1092fd39a41e82686ff", +- "https://bcr.bazel.build/modules/rules_pkg/1.0.1/source.json": "bd82e5d7b9ce2d31e380dd9f50c111d678c3bdaca190cb76b0e1c71b05e1ba8a", +-@@ -268,7 +273,8 @@ +- "https://bcr.bazel.build/modules/rules_python/1.6.0/MODULE.bazel": "7e04ad8f8d5bea40451cf80b1bd8262552aa73f841415d20db96b7241bd027d8", +- "https://bcr.bazel.build/modules/rules_python/1.7.0/MODULE.bazel": "d01f995ecd137abf30238ad9ce97f8fc3ac57289c8b24bd0bf53324d937a14f8", +- "https://bcr.bazel.build/modules/rules_python/1.8.0/MODULE.bazel": "c151c025dbcc93d8f62ab68ecc313c9176a868a0e6386981bf2a12aec77cbe7b", +-- "https://bcr.bazel.build/modules/rules_python/1.8.0/source.json": "356397eed5b46971d8c585c92098d70495078a80bf18bebcb4209f44b495f3e6", +-+ "https://bcr.bazel.build/modules/rules_python/1.8.4/MODULE.bazel": "33e3971e66161a3e955f7a0d411a8d1f291c4ce4c561851512466f3c77ff8ece", +-+ "https://bcr.bazel.build/modules/rules_python/1.8.4/source.json": "9fbc0e57bae52cddcc3831d668bce87a47e0c655104a85098d4459dd9a3b0a10", +- "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/MODULE.bazel": "d44fec647d0aeb67b9f3b980cf68ba634976f3ae7ccd6c07d790b59b87a4f251", +- "https://bcr.bazel.build/modules/rules_robolectric/4.14.1.2/source.json": "37c10335f2361c337c5c1f34ed36d2da70534c23088062b33a8bdaab68aa9dea", +- "https://bcr.bazel.build/modules/rules_rust/0.69.0/MODULE.bazel": "4326fec48f2fef0d514de46346f7f77e200c82936dd08b91c9ef039fbdad5c10", +-@@ -348,9 +354,142 @@ +- ] +- } +- }, +-+ "@@emsdk+//:emscripten_cache.bzl%emscripten_cache": { +-+ "general": { +-+ "bzlTransitiveDigest": "GMscy7c4sDbvbf9dMSrtUvyJJBuxYpJBT/Lg+2ob6dk=", +-+ "usagesDigest": "Id/C4z1d3MUlKZgmBOQiLi2E7NQZTGsd4WV2wSMzmo4=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": {}, +-+ "generatedRepoSpecs": { +-+ "emscripten_cache": { +-+ "repoRuleId": "@@emsdk+//:emscripten_cache.bzl%_emscripten_cache_repository", +-+ "attributes": { +-+ "configuration": [], +-+ "targets": [], +-+ "prebuilt_cache_url": "", +-+ "prebuilt_cache_sha256": "", +-+ "prebuilt_cache_strip_prefix": "" +-+ } +-+ } +-+ }, +-+ "recordedRepoMappingEntries": [] +-+ } +-+ }, +-+ "@@emsdk+//:emscripten_deps.bzl%emscripten_deps": { +-+ "general": { +-+ "bzlTransitiveDigest": "ZT33Pf8H8gJ/X4gT3oXVPzuAO0H2y1HtHMPOfILGV0M=", +-+ "usagesDigest": "4SlUap0Npa9PDUrLoi0uZ7CRDp9h/YbWi0UuidttuWc=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": {}, +-+ "generatedRepoSpecs": { +-+ "emscripten_bin_linux": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", +-+ "sha256": "d574428df9ecf00790e28636bdc47027432737c31621b18cdb418123afda4ac1", +-+ "strip_prefix": "install", +-+ "type": "tar.xz", +-+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" +-+ } +-+ }, +-+ "emscripten_bin_linux_arm64": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", +-+ "sha256": "d74803ef563511b9cc1e5cde5016f06d161ffd2b6223135a8aeeef44194594e7", +-+ "strip_prefix": "install", +-+ "type": "tar.xz", +-+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/linux/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" +-+ } +-+ }, +-+ "emscripten_bin_mac": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", +-+ "sha256": "356f36ba04a54edb029c658dd1b547c5c8a8f3c166b09654c1efcb9cf7bf8a57", +-+ "strip_prefix": "install", +-+ "type": "tar.xz", +-+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.tar.xz" +-+ } +-+ }, +-+ "emscripten_bin_mac_arm64": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang\",\n \"bin/clang++\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang\",\n \"bin/llvm-ar\",\n \"bin/llvm-dwarfdump\",\n \"bin/llvm-nm\",\n \"bin/llvm-objcopy\",\n \"bin/wasm-ctor-eval\",\n \"bin/wasm-emscripten-finalize\",\n \"bin/wasm-ld\",\n \"bin/wasm-metadce\",\n \"bin/wasm-opt\",\n \"bin/wasm-split\",\n \"bin/wasm2js\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", +-+ "sha256": "ded3bb783e7aa3dda576955dd0aa3a71dd21789e42befb63ee14f7d9f9b6aa32", +-+ "strip_prefix": "install", +-+ "type": "tar.xz", +-+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/mac/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries-arm64.tar.xz" +-+ } +-+ }, +-+ "emscripten_bin_win": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "build_file_content": "\npackage(default_visibility = ['//visibility:public'])\n\nfilegroup(\n name = \"all\",\n srcs = glob([\"**\"]),\n)\n\nfilegroup(\n name = \"includes\",\n srcs = glob([\n \"emscripten/cache/sysroot/include/c++/v1/**\",\n \"emscripten/cache/sysroot/include/compat/**\",\n \"emscripten/cache/sysroot/include/**\",\n \"lib/clang/**/include/**\",\n ]),\n)\n\nfilegroup(\n name = \"builtin_cache\",\n srcs = glob([\n \"emscripten/cache/**\",\n ]),\n)\n\nfilegroup(\n name = \"emcc_common\",\n srcs = [\n \"emscripten/emcc.py\",\n \"emscripten/embuilder.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/cache/sysroot_install.stamp\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/third_party/**\",\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"compiler_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/clang++.exe\",\n \":emcc_common\",\n \":includes\",\n ],\n)\n\nfilegroup(\n name = \"linker_files\",\n srcs = [\n \"bin/clang.exe\",\n \"bin/llvm-ar.exe\",\n \"bin/llvm-dwarfdump.exe\",\n \"bin/llvm-nm.exe\",\n \"bin/llvm-objcopy.exe\",\n \"bin/wasm-ctor-eval.exe\",\n \"bin/wasm-emscripten-finalize.exe\",\n \"bin/wasm-ld.exe\",\n \"bin/wasm-metadce.exe\",\n \"bin/wasm-opt.exe\",\n \"bin/wasm-split.exe\",\n \"bin/wasm2js.exe\",\n \":emcc_common\",\n ] + glob(\n include = [\n \"emscripten/cache/sysroot/lib/**\",\n \"emscripten/node_modules/**\",\n \"emscripten/src/**\",\n ],\n ),\n)\n\nfilegroup(\n name = \"ar_files\",\n srcs = [\n \"bin/llvm-ar.exe\",\n \"emscripten/emar.py\",\n \"emscripten/emscripten-version.txt\",\n \"emscripten/src/settings.js\",\n \"emscripten/src/settings_internal.js\",\n ] + glob(\n include = [\n \"emscripten/tools/**\",\n ],\n exclude = [\n \"**/__pycache__/**\",\n ],\n ),\n)\n", +-+ "sha256": "e5f9250a9cf4ff6ed16d57d6b5e177c844067381d31cd0c1a607c1ee1d2ba088", +-+ "strip_prefix": "install", +-+ "type": "zip", +-+ "url": "https://storage.googleapis.com/webassembly/emscripten-releases-builds/win/004876f1984e18a9eb0736c5ca417ac86d386fb8/wasm-binaries.zip" +-+ } +-+ } +-+ }, +-+ "recordedRepoMappingEntries": [ +-+ [ +-+ "emsdk+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "emsdk+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "cc_compatibility_proxy", +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", +-+ "rules_cc", +-+ "rules_cc+" +-+ ] +-+ ] +-+ } +-+ }, +-+ "@@protobuf+//python/dist:system_python.bzl%system_python_extension": { +-+ "general": { +-+ "bzlTransitiveDigest": "qh0n9IrXU/xS94wxKQrG1J63zrLkA1Wy2Y3BQxptPcI=", +-+ "usagesDigest": "tCi55FyqtOJ2jXh9vcjrHCl4ov3kpWiwKl103nA9BOI=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": {}, +-+ "generatedRepoSpecs": { +-+ "system_python": { +-+ "repoRuleId": "@@protobuf+//python/dist:system_python.bzl%system_python", +-+ "attributes": { +-+ "minimum_python_version": "3.9" +-+ } +-+ } +-+ }, +-+ "recordedRepoMappingEntries": [] +-+ } +-+ }, +- "@@pybind11_bazel+//:internal_configure.bzl%internal_configure_extension": { +- "general": { +-- "bzlTransitiveDigest": "NFQjcZF+fAvf5fDH+pqsx4JrfzP9PuHBz6S6ZutIbnw=", +-+ "bzlTransitiveDigest": "7zBsfo5dyMqKT23rXrvWqJMx0AugwL6NyirkmvzKcqU=", +- "usagesDigest": "D1r3lfzMuUBFxgG8V6o0bQTLMk3GkaGOaPzw53wrwyw=", +- "recordedFileInputs": { +- "@@pybind11_bazel+//MODULE.bazel": "e6f4c20442eaa7c90d7190d8dc539d0ab422f95c65a57cc59562170c58ae3d34" +-@@ -380,8 +519,8 @@ +- }, +- "@@rules_android+//bzlmod_extensions:apksig.bzl%apksig_extension": { +- "general": { +-- "bzlTransitiveDigest": "By9qVNN7G4oL1vYOJXye7Dp/CbR2ar9oxAW8WXAVcVw=", +-- "usagesDigest": "xq6OVkELeJvOgYo3oY/sUBsGFbcqdV+9BYiNgSPV/po=", +-+ "bzlTransitiveDigest": "15xx/lo4VYL9KdLW0Cc94ebALMF07+XZH8dZcVU8/LI=", +-+ "usagesDigest": "S8lLnnZxdeYUYq3kIGhVMk0wQ9Fd6elmCskvn+SL6iw=", +- "recordedFileInputs": {}, +- "recordedDirentsInputs": {}, +- "envVariables": {}, +-@@ -389,7 +528,10 @@ +- "apksig": { +- "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +- "attributes": { +-- "url": "https://android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz", +-+ "urls": [ +-+ "https://mirror.bazel.build/android.googlesource.com/platform/tools/apksig/+archive/24e3075e68ebe17c0b529bb24bfda819db5e2f3b.tar.gz" +-+ ], +-+ "sha256": "12e44fdbd219c5e1cc62099c2a01d775957603d2d4f693f8285f9d95d9a04e77", +- "build_file": "@@rules_android+//bzlmod_extensions:apksig.BUILD" +- } +- } +-@@ -405,8 +547,8 @@ +- }, +- "@@rules_android+//bzlmod_extensions:com_android_dex.bzl%com_android_dex_extension": { +- "general": { +-- "bzlTransitiveDigest": "rvWbJQc8jInfIAaXIMhSOqUlwM9HVeLey6q0ISvg08Y=", +-- "usagesDigest": "toF8IFMu98H/VU2p1sfVC5fVXVYJunpbbmtM6tOsQXY=", +-+ "bzlTransitiveDigest": "K0jbWcRwfM8njdIXNRjRvdApKmBfKeFLScDH+5LSSE0=", +-+ "usagesDigest": "0hluQmaWiWak6sVMP5L4wXhNyIwv9fw0y5JJ8lnPb1c=", +- "recordedFileInputs": {}, +- "recordedDirentsInputs": {}, +- "envVariables": {}, +-@@ -414,8 +556,11 @@ +- "com_android_dex": { +- "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +- "attributes": { +-- "url": "https://android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz", +-- "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD" +-+ "urls": [ +-+ "https://mirror.bazel.build/android.googlesource.com/platform/dalvik/+archive/5a81c499a569731e2395f7c8d13c0e0d4e17a2b6.tar.gz" +-+ ], +-+ "build_file": "@@rules_android+//bzlmod_extensions:com_android_dex.BUILD", +-+ "sha256": "86b4848c038bf687fadc812239cb01fb8d1d15cef3125b480a0448360992b95d" +- } +- } +- }, +-@@ -444,10 +589,144 @@ +- "recordedRepoMappingEntries": [] +- } +- }, +-+ "@@rules_nodejs+//nodejs:extensions.bzl%node": { +-+ "general": { +-+ "bzlTransitiveDigest": "4pUxCNc22K4I+6+4Nxu52Hur12tFRfa1JMsN5mdDv60=", +-+ "usagesDigest": "dqOjZvNvw6/DVBPAiKrXJNA0Tx4GT4Vj/VdUyGMpDL8=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": {}, +-+ "generatedRepoSpecs": { +-+ "nodejs_linux_amd64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "linux_amd64" +-+ } +-+ }, +-+ "nodejs_linux_arm64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "linux_arm64" +-+ } +-+ }, +-+ "nodejs_linux_s390x": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "linux_s390x" +-+ } +-+ }, +-+ "nodejs_linux_ppc64le": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "linux_ppc64le" +-+ } +-+ }, +-+ "nodejs_darwin_amd64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "darwin_amd64" +-+ } +-+ }, +-+ "nodejs_darwin_arm64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "darwin_arm64" +-+ } +-+ }, +-+ "nodejs_windows_amd64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "windows_amd64" +-+ } +-+ }, +-+ "nodejs_windows_arm64": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs:repositories.bzl%_nodejs_repositories", +-+ "attributes": { +-+ "node_download_auth": {}, +-+ "node_repositories": {}, +-+ "node_urls": [ +-+ "https://nodejs.org/dist/v{version}/{filename}" +-+ ], +-+ "node_version": "20.18.0", +-+ "include_headers": false, +-+ "platform": "windows_arm64" +-+ } +-+ }, +-+ "nodejs": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", +-+ "attributes": { +-+ "user_node_repository_name": "nodejs" +-+ } +-+ }, +-+ "nodejs_host": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_repo_host_os_alias.bzl%nodejs_repo_host_os_alias", +-+ "attributes": { +-+ "user_node_repository_name": "nodejs" +-+ } +-+ }, +-+ "nodejs_toolchains": { +-+ "repoRuleId": "@@rules_nodejs+//nodejs/private:nodejs_toolchains_repo.bzl%nodejs_toolchains_repo", +-+ "attributes": { +-+ "user_node_repository_name": "nodejs" +-+ } +-+ } +-+ }, +-+ "recordedRepoMappingEntries": [] +-+ } +-+ }, +- "@@rules_python+//python/extensions:config.bzl%config": { +- "general": { +- "bzlTransitiveDigest": "EcMcbtKZvYmd5Mi1Fpg4EeBBztLHEE5tjO5tLDBYDuU=", +-- "usagesDigest": "EocbSr4I3/Shk4QaFool8b8navUiUFqgzF9bOaaYfFk=", +-+ "usagesDigest": "p2al+dDKI5UlCyNvheMVynbWSGbdiji/jMz53fMNfJA=", +- "recordedFileInputs": {}, +- "recordedDirentsInputs": {}, +- "envVariables": {}, +-@@ -682,7 +961,7 @@ +- "@@rules_python+//python/uv:uv.bzl%uv": { +- "general": { +- "bzlTransitiveDigest": "ijW9KS7qsIY+yBVvJ+Nr1mzwQox09j13DnE3iIwaeTM=", +-- "usagesDigest": "yXvWfXAzpBeW71mWgwU3AqAzXX/dFACnx12eYvBsJ8w=", +-+ "usagesDigest": "/HRt5Hw/vpDr9CDrKEPjeDIjxo4307VLxMu8BNAEDWA=", +- "recordedFileInputs": {}, +- "recordedDirentsInputs": {}, +- "envVariables": {}, +-@@ -719,6 +998,533 @@ +- ] +- ] +- } +-+ }, +-+ "@@rules_rust+//crate_universe:extension.bzl%crate": { +-+ "general": { +-+ "bzlTransitiveDigest": "VVbU93QvGxFMzb9BcpYTYyyYDpj10Ya6Zm5RH1JEUhw=", +-+ "usagesDigest": "EuFUqVKVHF263jHTWOHXs4tFACdRNVOhwpoytdk19bs=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": { +-+ "CARGO_BAZEL_DEBUG": null, +-+ "CARGO_BAZEL_GENERATOR_SHA256": null, +-+ "CARGO_BAZEL_GENERATOR_URL": null, +-+ "CARGO_BAZEL_ISOLATED": null, +-+ "CARGO_BAZEL_REPIN": null, +-+ "CARGO_BAZEL_REPIN_ONLY": null, +-+ "CARGO_BAZEL_TIMEOUT": null, +-+ "REPIN": null +-+ }, +-+ "generatedRepoSpecs": { +-+ "crates": { +-+ "repoRuleId": "@@rules_rust+//crate_universe:extensions.bzl%_generate_repo", +-+ "attributes": { +-+ "contents": { +-+ "BUILD.bazel": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\npackage(default_visibility = [\"//visibility:public\"])\n\nexports_files(\n [\n \"cargo-bazel.json\",\n \"crates.bzl\",\n \"defs.bzl\",\n ] + glob(\n allow_empty = True,\n include = [\"*.bazel\"],\n ),\n)\n\nfilegroup(\n name = \"srcs\",\n srcs = glob(\n allow_empty = True,\n include = [\n \"*.bazel\",\n \"*.bzl\",\n ],\n ),\n)\n\n# Workspace Member Dependencies\nalias(\n name = \"googletest-0.14.3\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"googletest\",\n actual = \"@crates__googletest-0.14.3//:googletest\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme-0.3.37\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"linkme\",\n actual = \"@crates__linkme-0.3.37//:linkme\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste-1.0.15\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"paste\",\n actual = \"@crates__paste-1.0.15//:paste\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote-1.0.47\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"quote\",\n actual = \"@crates__quote-1.0.47//:quote\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn-3.0.3\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n\nalias(\n name = \"syn\",\n actual = \"@crates__syn-3.0.3//:syn\",\n tags = [\"manual\"],\n)\n", +-+ "alias_rules.bzl": "\"\"\"Alias that transitions its target to `compilation_mode=opt`. Use `transition_alias=\"opt\"` to enable.\"\"\"\n\nload(\"@rules_cc//cc:defs.bzl\", \"CcInfo\")\nload(\"@rules_rust//rust:rust_common.bzl\", \"COMMON_PROVIDERS\")\n\ndef _transition_alias_impl(ctx):\n # `ctx.attr.actual` is a list of 1 item due to the transition\n providers = [ctx.attr.actual[0][provider] for provider in COMMON_PROVIDERS]\n if CcInfo in ctx.attr.actual[0]:\n providers.append(ctx.attr.actual[0][CcInfo])\n return providers\n\ndef _change_compilation_mode(compilation_mode):\n def _change_compilation_mode_impl(_settings, _attr):\n return {\n \"//command_line_option:compilation_mode\": compilation_mode,\n }\n\n return transition(\n implementation = _change_compilation_mode_impl,\n inputs = [],\n outputs = [\n \"//command_line_option:compilation_mode\",\n ],\n )\n\ndef _transition_alias_rule(compilation_mode):\n return rule(\n implementation = _transition_alias_impl,\n provides = COMMON_PROVIDERS,\n attrs = {\n \"actual\": attr.label(\n mandatory = True,\n doc = \"`rust_library()` target to transition to `compilation_mode=opt`.\",\n providers = COMMON_PROVIDERS,\n cfg = _change_compilation_mode(compilation_mode),\n ),\n \"_allowlist_function_transition\": attr.label(\n default = \"@bazel_tools//tools/allowlists/function_transition_allowlist\",\n ),\n },\n doc = \"Transitions a Rust library crate to the `compilation_mode=opt`.\",\n )\n\ntransition_alias_dbg = _transition_alias_rule(\"dbg\")\ntransition_alias_fastbuild = _transition_alias_rule(\"fastbuild\")\ntransition_alias_opt = _transition_alias_rule(\"opt\")\n", +-+ "defs.bzl": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\"\"\"\n# `crates_repository` API\n\n- [aliases](#aliases)\n- [crate_deps](#crate_deps)\n- [all_crate_deps](#all_crate_deps)\n- [crate_repositories](#crate_repositories)\n\n\"\"\"\n\nload(\"@bazel_tools//tools/build_defs/repo:git.bzl\", \"new_git_repository\")\nload(\"@bazel_tools//tools/build_defs/repo:http.bzl\", \"http_archive\")\nload(\"@bazel_tools//tools/build_defs/repo:utils.bzl\", \"maybe\")\nload(\"@bazel_skylib//lib:selects.bzl\", \"selects\")\nload(\"@rules_rust//crate_universe/private:local_crate_mirror.bzl\", \"local_crate_mirror\")\n\n###############################################################################\n# MACROS API\n###############################################################################\n\n# An identifier that represent common dependencies (unconditional).\n_COMMON_CONDITION = \"\"\n\ndef _flatten_dependency_maps(all_dependency_maps):\n \"\"\"Flatten a list of dependency maps into one dictionary.\n\n Dependency maps have the following structure:\n\n ```python\n DEPENDENCIES_MAP = {\n # The first key in the map is a Bazel package\n # name of the workspace this file is defined in.\n \"workspace_member_package\": {\n\n # Not all dependencies are supported for all platforms.\n # the condition key is the condition required to be true\n # on the host platform.\n \"condition\": {\n\n # An alias to a crate target. # The label of the crate target the\n # Aliases are only crate names. # package name refers to.\n \"package_name\": \"@full//:label\",\n }\n }\n }\n ```\n\n Args:\n all_dependency_maps (list): A list of dicts as described above\n\n Returns:\n dict: A dictionary as described above\n \"\"\"\n dependencies = {}\n\n for workspace_deps_map in all_dependency_maps:\n for pkg_name, conditional_deps_map in workspace_deps_map.items():\n if pkg_name not in dependencies:\n non_frozen_map = dict()\n for key, values in conditional_deps_map.items():\n non_frozen_map.update({key: dict(values.items())})\n dependencies.setdefault(pkg_name, non_frozen_map)\n continue\n\n for condition, deps_map in conditional_deps_map.items():\n # If the condition has not been recorded, do so and continue\n if condition not in dependencies[pkg_name]:\n dependencies[pkg_name].setdefault(condition, dict(deps_map.items()))\n continue\n\n # Alert on any miss-matched dependencies\n inconsistent_entries = []\n for crate_name, crate_label in deps_map.items():\n existing = dependencies[pkg_name][condition].get(crate_name)\n if existing and existing != crate_label:\n inconsistent_entries.append((crate_name, existing, crate_label))\n dependencies[pkg_name][condition].update({crate_name: crate_label})\n\n return dependencies\n\ndef crate_deps(deps, package_name = None):\n \"\"\"Finds the fully qualified label of the requested crates for the package where this macro is called.\n\n Args:\n deps (list): The desired list of crate targets.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()`.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if not deps:\n return []\n\n if package_name == None:\n package_name = native.package_name()\n\n # Join both sets of dependencies\n dependencies = _flatten_dependency_maps([\n _NORMAL_DEPENDENCIES,\n _NORMAL_DEV_DEPENDENCIES,\n _PROC_MACRO_DEPENDENCIES,\n _PROC_MACRO_DEV_DEPENDENCIES,\n _BUILD_DEPENDENCIES,\n _BUILD_PROC_MACRO_DEPENDENCIES,\n ]).pop(package_name, {})\n\n # Combine all conditional packages so we can easily index over a flat list\n # TODO: Perhaps this should actually return select statements and maintain\n # the conditionals of the dependencies\n flat_deps = {}\n for deps_set in dependencies.values():\n for crate_name, crate_label in deps_set.items():\n flat_deps.update({crate_name: crate_label})\n\n missing_crates = []\n crate_targets = []\n for crate_target in deps:\n if crate_target not in flat_deps:\n missing_crates.append(crate_target)\n else:\n crate_targets.append(flat_deps[crate_target])\n\n if missing_crates:\n fail(\"Could not find crates `{}` among dependencies of `{}`. Available dependencies were `{}`\".format(\n missing_crates,\n package_name,\n dependencies,\n ))\n\n return crate_targets\n\ndef all_crate_deps(\n normal = False, \n normal_dev = False, \n proc_macro = False, \n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Finds the fully qualified label of all requested direct crate dependencies \\\n for the package where this macro is called.\n\n If no parameters are set, all normal dependencies are returned. Setting any one flag will\n otherwise impact the contents of the returned list.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list.\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n list: A list of labels to generated rust targets (str)\n \"\"\"\n\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_dependency_maps = []\n if normal:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n if normal_dev:\n all_dependency_maps.append(_NORMAL_DEV_DEPENDENCIES)\n if proc_macro:\n all_dependency_maps.append(_PROC_MACRO_DEPENDENCIES)\n if proc_macro_dev:\n all_dependency_maps.append(_PROC_MACRO_DEV_DEPENDENCIES)\n if build:\n all_dependency_maps.append(_BUILD_DEPENDENCIES)\n if build_proc_macro:\n all_dependency_maps.append(_BUILD_PROC_MACRO_DEPENDENCIES)\n\n # Default to always using normal dependencies\n if not all_dependency_maps:\n all_dependency_maps.append(_NORMAL_DEPENDENCIES)\n\n dependencies = _flatten_dependency_maps(all_dependency_maps).pop(package_name, None)\n\n if not dependencies:\n if dependencies == None:\n fail(\"Tried to get all_crate_deps for package \" + package_name + \" but that package had no Cargo.toml file\")\n else:\n return []\n\n crate_deps = list(dependencies.pop(_COMMON_CONDITION, {}).values())\n for condition, deps in dependencies.items():\n crate_deps += selects.with_or({\n tuple(_CONDITIONS[condition]): deps.values(),\n \"//conditions:default\": [],\n })\n\n return crate_deps\n\ndef aliases(\n normal = False,\n normal_dev = False,\n proc_macro = False,\n proc_macro_dev = False,\n build = False,\n build_proc_macro = False,\n package_name = None):\n \"\"\"Produces a map of Crate alias names to their original label\n\n If no dependency kinds are specified, `normal` and `proc_macro` are used by default.\n Setting any one flag will otherwise determine the contents of the returned dict.\n\n Args:\n normal (bool, optional): If True, normal dependencies are included in the\n output list.\n normal_dev (bool, optional): If True, normal dev dependencies will be\n included in the output list..\n proc_macro (bool, optional): If True, proc_macro dependencies are included\n in the output list.\n proc_macro_dev (bool, optional): If True, dev proc_macro dependencies are\n included in the output list.\n build (bool, optional): If True, build dependencies are included\n in the output list.\n build_proc_macro (bool, optional): If True, build proc_macro dependencies are\n included in the output list.\n package_name (str, optional): The package name of the set of dependencies to look up.\n Defaults to `native.package_name()` when unset.\n\n Returns:\n dict: The aliases of all associated packages\n \"\"\"\n if package_name == None:\n package_name = native.package_name()\n\n # Determine the relevant maps to use\n all_aliases_maps = []\n if normal:\n all_aliases_maps.append(_NORMAL_ALIASES)\n if normal_dev:\n all_aliases_maps.append(_NORMAL_DEV_ALIASES)\n if proc_macro:\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n if proc_macro_dev:\n all_aliases_maps.append(_PROC_MACRO_DEV_ALIASES)\n if build:\n all_aliases_maps.append(_BUILD_ALIASES)\n if build_proc_macro:\n all_aliases_maps.append(_BUILD_PROC_MACRO_ALIASES)\n\n # Default to always using normal aliases\n if not all_aliases_maps:\n all_aliases_maps.append(_NORMAL_ALIASES)\n all_aliases_maps.append(_PROC_MACRO_ALIASES)\n\n aliases = _flatten_dependency_maps(all_aliases_maps).pop(package_name, None)\n\n if not aliases:\n return dict()\n\n common_items = aliases.pop(_COMMON_CONDITION, {}).items()\n\n # If there are only common items in the dictionary, immediately return them\n if not len(aliases.keys()) == 1:\n return dict(common_items)\n\n # Build a single select statement where each conditional has accounted for the\n # common set of aliases.\n crate_aliases = {\"//conditions:default\": dict(common_items)}\n for condition, deps in aliases.items():\n condition_triples = _CONDITIONS[condition]\n for triple in condition_triples:\n if triple in crate_aliases:\n crate_aliases[triple].update(deps)\n else:\n crate_aliases.update({triple: dict(deps.items() + common_items)})\n\n return select(crate_aliases)\n\n###############################################################################\n# WORKSPACE MEMBER DEPS AND ALIASES\n###############################################################################\n\n_NORMAL_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"googletest\": Label(\"@crates//:googletest-0.14.3\"),\n \"linkme\": Label(\"@crates//:linkme-0.3.37\"),\n \"quote\": Label(\"@crates//:quote-1.0.47\"),\n \"syn\": Label(\"@crates//:syn-3.0.3\"),\n },\n },\n}\n\n\n_NORMAL_ALIASES = {\n \"\": {\n _COMMON_CONDITION: {\n },\n },\n}\n\n\n_NORMAL_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_NORMAL_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n _COMMON_CONDITION: {\n \"paste\": Label(\"@crates//:paste-1.0.15\"),\n },\n },\n}\n\n\n_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_PROC_MACRO_DEV_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_ALIASES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_DEPENDENCIES = {\n \"\": {\n },\n}\n\n\n_BUILD_PROC_MACRO_ALIASES = {\n \"\": {\n },\n}\n\n\n_CONDITIONS = {\n \"aarch64-apple-darwin\": [\"@rules_rust//rust/platform:aarch64-apple-darwin\"],\n \"aarch64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\"],\n \"wasm32-unknown-unknown\": [\"@rules_rust//rust/platform:wasm32-unknown-unknown\"],\n \"wasm32-wasip1\": [\"@rules_rust//rust/platform:wasm32-wasip1\"],\n \"x86_64-pc-windows-msvc\": [\"@rules_rust//rust/platform:x86_64-pc-windows-msvc\"],\n \"x86_64-unknown-linux-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\"],\n \"x86_64-unknown-nixos-gnu\": [\"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\"],\n}\n\n###############################################################################\n\ndef crate_repositories():\n \"\"\"A macro for defining repositories for all generated crates.\n\n Returns:\n A list of repos visible to the module through the module extension.\n \"\"\"\n maybe(\n http_archive,\n name = \"crates__aho-corasick-1.1.5\",\n sha256 = \"c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/aho-corasick/1.1.5/download\"],\n strip_prefix = \"aho-corasick-1.1.5\",\n build_file = Label(\"@crates//crates:BUILD.aho-corasick-1.1.5.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__autocfg-1.5.1\",\n sha256 = \"f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/autocfg/1.5.1/download\"],\n strip_prefix = \"autocfg-1.5.1\",\n build_file = Label(\"@crates//crates:BUILD.autocfg-1.5.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest-0.14.3\",\n sha256 = \"f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest/0.14.3/download\"],\n strip_prefix = \"googletest-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__googletest_macro-0.14.3\",\n sha256 = \"2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/googletest_macro/0.14.3/download\"],\n strip_prefix = \"googletest_macro-0.14.3\",\n build_file = Label(\"@crates//crates:BUILD.googletest_macro-0.14.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-0.3.37\",\n sha256 = \"3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme/0.3.37/download\"],\n strip_prefix = \"linkme-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__linkme-impl-0.3.37\",\n sha256 = \"77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/linkme-impl/0.3.37/download\"],\n strip_prefix = \"linkme-impl-0.3.37\",\n build_file = Label(\"@crates//crates:BUILD.linkme-impl-0.3.37.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__memchr-2.8.3\",\n sha256 = \"cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/memchr/2.8.3/download\"],\n strip_prefix = \"memchr-2.8.3\",\n build_file = Label(\"@crates//crates:BUILD.memchr-2.8.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__num-traits-0.2.19\",\n sha256 = \"071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/num-traits/0.2.19/download\"],\n strip_prefix = \"num-traits-0.2.19\",\n build_file = Label(\"@crates//crates:BUILD.num-traits-0.2.19.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__paste-1.0.15\",\n sha256 = \"57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/paste/1.0.15/download\"],\n strip_prefix = \"paste-1.0.15\",\n build_file = Label(\"@crates//crates:BUILD.paste-1.0.15.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__proc-macro2-1.0.107\",\n sha256 = \"985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/proc-macro2/1.0.107/download\"],\n strip_prefix = \"proc-macro2-1.0.107\",\n build_file = Label(\"@crates//crates:BUILD.proc-macro2-1.0.107.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__quote-1.0.47\",\n sha256 = \"1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/quote/1.0.47/download\"],\n strip_prefix = \"quote-1.0.47\",\n build_file = Label(\"@crates//crates:BUILD.quote-1.0.47.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-1.13.1\",\n sha256 = \"f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex/1.13.1/download\"],\n strip_prefix = \"regex-1.13.1\",\n build_file = Label(\"@crates//crates:BUILD.regex-1.13.1.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-automata-0.4.18\",\n sha256 = \"ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-automata/0.4.18/download\"],\n strip_prefix = \"regex-automata-0.4.18\",\n build_file = Label(\"@crates//crates:BUILD.regex-automata-0.4.18.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__regex-syntax-0.8.11\",\n sha256 = \"d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/regex-syntax/0.8.11/download\"],\n strip_prefix = \"regex-syntax-0.8.11\",\n build_file = Label(\"@crates//crates:BUILD.regex-syntax-0.8.11.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__rustversion-1.0.23\",\n sha256 = \"cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/rustversion/1.0.23/download\"],\n strip_prefix = \"rustversion-1.0.23\",\n build_file = Label(\"@crates//crates:BUILD.rustversion-1.0.23.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-2.0.119\",\n sha256 = \"872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/2.0.119/download\"],\n strip_prefix = \"syn-2.0.119\",\n build_file = Label(\"@crates//crates:BUILD.syn-2.0.119.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__syn-3.0.3\",\n sha256 = \"53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/syn/3.0.3/download\"],\n strip_prefix = \"syn-3.0.3\",\n build_file = Label(\"@crates//crates:BUILD.syn-3.0.3.bazel\"),\n )\n\n maybe(\n http_archive,\n name = \"crates__unicode-ident-1.0.24\",\n sha256 = \"e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75\",\n type = \"tar.gz\",\n urls = [\"https://static.crates.io/crates/unicode-ident/1.0.24/download\"],\n strip_prefix = \"unicode-ident-1.0.24\",\n build_file = Label(\"@crates//crates:BUILD.unicode-ident-1.0.24.bazel\"),\n )\n\n return [\n struct(repo=\"crates__googletest-0.14.3\", is_dev_dep = False),\n struct(repo=\"crates__linkme-0.3.37\", is_dev_dep = False),\n struct(repo=\"crates__paste-1.0.15\", is_dev_dep = False),\n struct(repo=\"crates__quote-1.0.47\", is_dev_dep = False),\n struct(repo=\"crates__syn-3.0.3\", is_dev_dep = False),\n ]\n" +-+ } +-+ } +-+ }, +-+ "crates__aho-corasick-1.1.5": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "c982642fa9e8606056828ee9a8505737230110bb1099153c79efe865c59d12ba", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/aho-corasick/1.1.5/download" +-+ ], +-+ "strip_prefix": "aho-corasick-1.1.5", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"aho_corasick\",\n deps = [\n \"@crates__memchr-2.8.3//:memchr\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"perf-literal\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=aho-corasick\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.1.5\",\n)\n" +-+ } +-+ }, +-+ "crates__autocfg-1.5.1": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "f2032f911046de80f0a198e0901378627c33f59ea0ac00e363d481118bd70a53", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/autocfg/1.5.1/download" +-+ ], +-+ "strip_prefix": "autocfg-1.5.1", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"autocfg\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2015\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=autocfg\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.5.1\",\n)\n" +-+ } +-+ }, +-+ "crates__googletest-0.14.3": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "f6b5e2f2b556b7b90297a5a35c8267dd43a537923d2b329beefdba2b4ec19d94", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/googletest/0.14.3/download" +-+ ], +-+ "strip_prefix": "googletest-0.14.3", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"googletest\",\n deps = [\n \"@crates__num-traits-0.2.19//:num_traits\",\n \"@crates__regex-1.13.1//:regex\",\n ],\n proc_macro_deps = [\n \"@crates__googletest_macro-0.14.3//:googletest_macro\",\n \"@crates__rustversion-1.0.23//:rustversion\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" +-+ } +-+ }, +-+ "crates__googletest_macro-0.14.3": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "2ae6abc96141edd26bf5aeec0f119c129c44de3ced09e5073711a02cb74725d0", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/googletest_macro/0.14.3/download" +-+ ], +-+ "strip_prefix": "googletest_macro-0.14.3", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"googletest_macro\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-2.0.119//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=googletest_macro\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.14.3\",\n)\n" +-+ } +-+ }, +-+ "crates__linkme-0.3.37": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "3045e122bd98aef8ec3ad58ce84f0791f64e70163d1a02710af4aa11a4d54cc5", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/linkme/0.3.37/download" +-+ ], +-+ "strip_prefix": "linkme-0.3.37", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"linkme\",\n deps = [\n \"@crates__linkme-0.3.37//:build_script_build\",\n ],\n proc_macro_deps = [\n \"@crates__linkme-impl-0.3.37//:linkme_impl\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__linkme-impl-0.3.37": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "77060ebe535362c3da75682cd17b0431017b6e7c5661e714fc69a7ad017d1301", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/linkme-impl/0.3.37/download" +-+ ], +-+ "strip_prefix": "linkme-impl-0.3.37", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"linkme_impl\",\n deps = [\n \"@crates__linkme-impl-0.3.37//:build_script_build\",\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__syn-3.0.3//:syn\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.3.37\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"linkme-impl\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=linkme-impl\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.3.37\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__memchr-2.8.3": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "cf8baf1c55e62ffcace7a9f06f4bd9cd3f0c4beb022d3b367256b91b87513d98", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/memchr/2.8.3/download" +-+ ], +-+ "strip_prefix": "memchr-2.8.3", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"memchr\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=memchr\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.8.3\",\n)\n" +-+ } +-+ }, +-+ "crates__num-traits-0.2.19": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "071dfc062690e90b734c0b2273ce72ad0ffa95f0c74596bc250dcfd960262841", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/num-traits/0.2.19/download" +-+ ], +-+ "strip_prefix": "num-traits-0.2.19", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"num_traits\",\n deps = [\n \"@crates__num-traits-0.2.19//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.2.19\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n deps = [\n \"@crates__autocfg-1.5.1//:autocfg\",\n ],\n edition = \"2021\",\n pkg_name = \"num-traits\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=num-traits\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"0.2.19\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__paste-1.0.15": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "57c0d7b74b563b49d38dae00a0c37d4d6de9b432382b2892f0574ddcae73fd0a", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/paste/1.0.15/download" +-+ ], +-+ "strip_prefix": "paste-1.0.15", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"paste\",\n deps = [\n \"@crates__paste-1.0.15//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.15\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"paste\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=paste\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.15\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__proc-macro2-1.0.107": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "985e7ec9bb745e6ce6535b544d84d6cd6f7ad8bd711c398938ae983b91a766d9", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/proc-macro2/1.0.107/download" +-+ ], +-+ "strip_prefix": "proc-macro2-1.0.107", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"proc_macro2\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:build_script_build\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.107\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"proc-macro\",\n ] + select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [\n \"default\", # aarch64-apple-darwin\n ],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [\n \"default\", # aarch64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [\n \"default\", # x86_64-pc-windows-msvc\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [\n \"default\", # x86_64-unknown-linux-gnu\n ],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [\n \"default\", # x86_64-unknown-nixos-gnu\n ],\n \"//conditions:default\": [],\n }),\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"proc-macro2\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=proc-macro2\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.107\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__quote-1.0.47": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "1fbf4db142a473a8d80c26bbf18454ed458bf8d26c8219c331daecfdbd079001", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/quote/1.0.47/download" +-+ ], +-+ "strip_prefix": "quote-1.0.47", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"quote\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.47\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"proc-macro\",\n ],\n crate_name = \"build_script_build\",\n crate_root = \"build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2021\",\n pkg_name = \"quote\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=quote\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.47\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__regex-1.13.1": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "f020237b6c8eed93db2e2cb53c00c60a8e1bc73da7d073199a1180401450218d", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/regex/1.13.1/download" +-+ ], +-+ "strip_prefix": "regex-1.13.1", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-automata-0.4.18//:regex_automata\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"perf\",\n \"perf-backtrack\",\n \"perf-cache\",\n \"perf-dfa\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-onepass\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.13.1\",\n)\n" +-+ } +-+ }, +-+ "crates__regex-automata-0.4.18": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "ad8553b9b26413251cbf30e620595c7a41b3887f03da04579c0e6b0d6a06b4b2", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/regex-automata/0.4.18/download" +-+ ], +-+ "strip_prefix": "regex-automata-0.4.18", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_automata\",\n deps = [\n \"@crates__aho-corasick-1.1.5//:aho_corasick\",\n \"@crates__memchr-2.8.3//:memchr\",\n \"@crates__regex-syntax-0.8.11//:regex_syntax\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"alloc\",\n \"dfa-onepass\",\n \"hybrid\",\n \"meta\",\n \"nfa-backtrack\",\n \"nfa-pikevm\",\n \"nfa-thompson\",\n \"perf-inline\",\n \"perf-literal\",\n \"perf-literal-multisubstring\",\n \"perf-literal-substring\",\n \"std\",\n \"syntax\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n \"unicode-word-boundary\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-automata\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.4.18\",\n)\n" +-+ } +-+ }, +-+ "crates__regex-syntax-0.8.11": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "d6f6ff9a378485b298a5286656da665ba74413d36db0979633275d2e708145d4", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/regex-syntax/0.8.11/download" +-+ ], +-+ "strip_prefix": "regex-syntax-0.8.11", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"regex_syntax\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"default\",\n \"std\",\n \"unicode\",\n \"unicode-age\",\n \"unicode-bool\",\n \"unicode-case\",\n \"unicode-gencat\",\n \"unicode-perl\",\n \"unicode-script\",\n \"unicode-segment\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=regex-syntax\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"0.8.11\",\n)\n" +-+ } +-+ }, +-+ "crates__rustversion-1.0.23": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "cf54715a573b99ac80df0bc206da022bcd442c974952c7b9720069370852e21f", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/rustversion/1.0.23/download" +-+ ], +-+ "strip_prefix": "rustversion-1.0.23", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\n \"@rules_rust//cargo:defs.bzl\",\n \"cargo_build_script\",\n \"cargo_toml_env_vars\",\n)\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_proc_macro\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_proc_macro(\n name = \"rustversion\",\n deps = [\n \"@crates__rustversion-1.0.23//:build_script_build\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2018\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.23\",\n)\n\ncargo_build_script(\n name = \"_bs\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \"**/*.rs\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_name = \"build_script_build\",\n crate_root = \"build/build.rs\",\n data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n edition = \"2018\",\n pkg_name = \"rustversion\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=rustversion\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n version = \"1.0.23\",\n visibility = [\"//visibility:private\"],\n)\n\nalias(\n name = \"build_script_build\",\n actual = \":_bs\",\n tags = [\"manual\"],\n)\n" +-+ } +-+ }, +-+ "crates__syn-2.0.119": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "872831b642d1a07999a962a351ed35b955ea2cfc8f3862091e2a240a84f17297", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/syn/2.0.119/download" +-+ ], +-+ "strip_prefix": "syn-2.0.119", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"extra-traits\",\n \"full\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"2.0.119\",\n)\n" +-+ } +-+ }, +-+ "crates__syn-3.0.3": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "53e9bae58849f64dfa4f5d5ae372c8341f7305f82a3868709269343628b659a3", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/syn/3.0.3/download" +-+ ], +-+ "strip_prefix": "syn-3.0.3", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"syn\",\n deps = [\n \"@crates__proc-macro2-1.0.107//:proc_macro2\",\n \"@crates__quote-1.0.47//:quote\",\n \"@crates__unicode-ident-1.0.24//:unicode_ident\",\n ],\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_features = [\n \"clone-impls\",\n \"default\",\n \"derive\",\n \"parsing\",\n \"printing\",\n \"proc-macro\",\n ],\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=syn\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"3.0.3\",\n)\n" +-+ } +-+ }, +-+ "crates__unicode-ident-1.0.24": { +-+ "repoRuleId": "@@bazel_tools//tools/build_defs/repo:http.bzl%http_archive", +-+ "attributes": { +-+ "patch_args": [], +-+ "patch_tool": "", +-+ "patches": [], +-+ "remote_patch_strip": 1, +-+ "sha256": "e6e4313cd5fcd3dad5cafa179702e2b244f760991f45397d14d4ebf38247da75", +-+ "type": "tar.gz", +-+ "urls": [ +-+ "https://static.crates.io/crates/unicode-ident/1.0.24/download" +-+ ], +-+ "strip_prefix": "unicode-ident-1.0.24", +-+ "build_file_content": "###############################################################################\n# @generated\n# DO NOT MODIFY: This file is auto-generated by a crate_universe tool. To \n# regenerate this file, run the following:\n#\n# bazel mod show_repo 'protobuf'\n###############################################################################\n\nload(\"@rules_rust//cargo:defs.bzl\", \"cargo_toml_env_vars\")\n\nload(\"@rules_rust//rust:defs.bzl\", \"rust_library\")\n\n# buildifier: disable=bzl-visibility\nload(\"@rules_rust//crate_universe/private:selects.bzl\", \"selects\")\n\npackage(default_visibility = [\"//visibility:public\"])\n\ncargo_toml_env_vars(\n name = \"cargo_toml_env_vars\",\n src = \"Cargo.toml\",\n)\n\nrust_library(\n name = \"unicode_ident\",\n compile_data = glob(\n allow_empty = True,\n include = [\"**\"],\n exclude = [\n \"**/* *\",\n \".tmp_git_root/**/*\",\n \"BUILD\",\n \"BUILD.bazel\",\n \"WORKSPACE\",\n \"WORKSPACE.bazel\",\n ],\n ),\n crate_root = \"src/lib.rs\",\n edition = \"2021\",\n rustc_env_files = [\n \":cargo_toml_env_vars\",\n ],\n rustc_flags = [\n \"--cap-lints=allow\",\n ],\n srcs = glob(\n allow_empty = True,\n include = [\"**/*.rs\"],\n ),\n tags = [\n \"cargo-bazel\",\n \"crate-name=unicode-ident\",\n \"manual\",\n \"noclippy\",\n \"norustfmt\",\n ],\n target_compatible_with = select({\n \"@rules_rust//rust/platform:aarch64-apple-darwin\": [],\n \"@rules_rust//rust/platform:aarch64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:wasm32-unknown-unknown\": [],\n \"@rules_rust//rust/platform:wasm32-wasip1\": [],\n \"@rules_rust//rust/platform:x86_64-pc-windows-msvc\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-linux-gnu\": [],\n \"@rules_rust//rust/platform:x86_64-unknown-nixos-gnu\": [],\n \"//conditions:default\": [\"@platforms//:incompatible\"],\n }),\n version = \"1.0.24\",\n)\n" +-+ } +-+ } +-+ }, +-+ "recordedRepoMappingEntries": [ +-+ [ +-+ "bazel_features+", +-+ "bazel_features_globals", +-+ "bazel_features++version_extension+bazel_features_globals" +-+ ], +-+ [ +-+ "bazel_features+", +-+ "bazel_features_version", +-+ "bazel_features++version_extension+bazel_features_version" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "cc_compatibility_proxy", +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_features", +-+ "bazel_features+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_skylib", +-+ "bazel_skylib+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "rules_rust", +-+ "rules_rust+" +-+ ] +-+ ] +-+ } +-+ }, +-+ "@@rules_rust+//crate_universe/private:internal_extensions.bzl%cu_nr": { +-+ "general": { +-+ "bzlTransitiveDigest": "GOOgbXFJQhO4daGipwnspaixIHp6AWTvXRBe2wMULd4=", +-+ "usagesDigest": "tG3p3Nb5XxC7vWY/bcKdb//g0HoAxpxxH3F5/jBVlk4=", +-+ "recordedFileInputs": {}, +-+ "recordedDirentsInputs": {}, +-+ "envVariables": {}, +-+ "generatedRepoSpecs": { +-+ "cargo_bazel_bootstrap": { +-+ "repoRuleId": "@@rules_rust+//cargo/private:cargo_bootstrap.bzl%cargo_bootstrap_repository", +-+ "attributes": { +-+ "srcs": [ +-+ "@@rules_rust+//crate_universe:src/api.rs", +-+ "@@rules_rust+//crate_universe:src/api/lockfile.rs", +-+ "@@rules_rust+//crate_universe:src/cli.rs", +-+ "@@rules_rust+//crate_universe:src/cli/generate.rs", +-+ "@@rules_rust+//crate_universe:src/cli/query.rs", +-+ "@@rules_rust+//crate_universe:src/cli/render.rs", +-+ "@@rules_rust+//crate_universe:src/cli/splice.rs", +-+ "@@rules_rust+//crate_universe:src/cli/vendor.rs", +-+ "@@rules_rust+//crate_universe:src/config.rs", +-+ "@@rules_rust+//crate_universe:src/context.rs", +-+ "@@rules_rust+//crate_universe:src/context/crate_context.rs", +-+ "@@rules_rust+//crate_universe:src/context/platforms.rs", +-+ "@@rules_rust+//crate_universe:src/lib.rs", +-+ "@@rules_rust+//crate_universe:src/lockfile.rs", +-+ "@@rules_rust+//crate_universe:src/main.rs", +-+ "@@rules_rust+//crate_universe:src/metadata.rs", +-+ "@@rules_rust+//crate_universe:src/metadata/cargo_bin.rs", +-+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_resolver.rs", +-+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.bat", +-+ "@@rules_rust+//crate_universe:src/metadata/cargo_tree_rustc_wrapper.sh", +-+ "@@rules_rust+//crate_universe:src/metadata/dependency.rs", +-+ "@@rules_rust+//crate_universe:src/metadata/metadata_annotation.rs", +-+ "@@rules_rust+//crate_universe:src/rendering.rs", +-+ "@@rules_rust+//crate_universe:src/rendering/template_engine.rs", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/module_bzl.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/header.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/aliases_map.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/deps_map.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_git.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/partials/module/repo_http.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/templates/vendor_module.j2", +-+ "@@rules_rust+//crate_universe:src/rendering/verbatim/alias_rules.bzl", +-+ "@@rules_rust+//crate_universe:src/select.rs", +-+ "@@rules_rust+//crate_universe:src/splicing.rs", +-+ "@@rules_rust+//crate_universe:src/splicing/cargo_config.rs", +-+ "@@rules_rust+//crate_universe:src/splicing/crate_index_lookup.rs", +-+ "@@rules_rust+//crate_universe:src/splicing/splicer.rs", +-+ "@@rules_rust+//crate_universe:src/test.rs", +-+ "@@rules_rust+//crate_universe:src/utils.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/glob.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/label.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/select.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/select_dict.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/select_list.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/select_scalar.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/select_set.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/serialize.rs", +-+ "@@rules_rust+//crate_universe:src/utils/starlark/target_compatible_with.rs", +-+ "@@rules_rust+//crate_universe:src/utils/symlink.rs", +-+ "@@rules_rust+//crate_universe:src/utils/target_triple.rs" +-+ ], +-+ "binary": "cargo-bazel", +-+ "cargo_lockfile": "@@rules_rust+//crate_universe:Cargo.lock", +-+ "cargo_toml": "@@rules_rust+//crate_universe:Cargo.toml", +-+ "version": "1.93.1", +-+ "timeout": 900, +-+ "rust_toolchain_cargo_template": "@rust_host_tools//:bin/{tool}", +-+ "rust_toolchain_rustc_template": "@rust_host_tools//:bin/{tool}", +-+ "compressed_windows_toolchain_names": false +-+ } +-+ } +-+ }, +-+ "moduleExtensionMetadata": { +-+ "explicitRootModuleDirectDeps": [ +-+ "cargo_bazel_bootstrap" +-+ ], +-+ "explicitRootModuleDirectDevDeps": [], +-+ "useAllRepos": "NO", +-+ "reproducible": false +-+ }, +-+ "recordedRepoMappingEntries": [ +-+ [ +-+ "bazel_features+", +-+ "bazel_features_globals", +-+ "bazel_features++version_extension+bazel_features_globals" +-+ ], +-+ [ +-+ "bazel_features+", +-+ "bazel_features_version", +-+ "bazel_features++version_extension+bazel_features_version" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "cc_compatibility_proxy", +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy" +-+ ], +-+ [ +-+ "rules_cc+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_cc++compatibility_proxy+cc_compatibility_proxy", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_features", +-+ "bazel_features+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_skylib", +-+ "bazel_skylib+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "bazel_tools", +-+ "bazel_tools" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "cui", +-+ "rules_rust++cu+cui" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "rrc", +-+ "rules_rust++i2+rrc" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "rules_cc", +-+ "rules_cc+" +-+ ], +-+ [ +-+ "rules_rust+", +-+ "rules_rust", +-+ "rules_rust+" +-+ ] +-+ ] +-+ } +- } +- }, +- "facts": { diff --git a/third_party/llvm/workspace.bzl b/third_party/llvm/workspace.bzl -index ad5fa7f9..7af3a857 100644 +index 7af3a857..08ae305f 100644 --- a/third_party/llvm/workspace.bzl +++ b/third_party/llvm/workspace.bzl @@ -19,8 +19,8 @@ load("//third_party:repo.bzl", "tf_http_archive") def repo(name): """Imports LLVM.""" -- LLVM_COMMIT = "fdb39b1112a6db0343de169b79820768bd16c7d9" -- LLVM_SHA256 = "15b2d88aa699b4c1d47b489fc0d8a4c389deee53d93c4a985e4c8dc251397da3" -+ LLVM_COMMIT = "ab547095ead5464dc024d66264d9b8a987f429f3" -+ LLVM_SHA256 = "c47662ae375e9523870b87006f27c7f35f686f1489c0cbdf571a583a66c973c9" +- LLVM_COMMIT = "ab547095ead5464dc024d66264d9b8a987f429f3" +- LLVM_SHA256 = "c47662ae375e9523870b87006f27c7f35f686f1489c0cbdf571a583a66c973c9" ++ LLVM_COMMIT = "cbc5a226cbf8d1d37d1ba8e55ce6973e8ef739cd" ++ LLVM_SHA256 = "1169d25c76701c535ecad40139ad48e4cec35d3e573308a2b5a2f63aaf2aa70a" tf_http_archive( name = name, -diff --git a/third_party/stablehlo/temporary.patch b/third_party/stablehlo/temporary.patch -index ad4e3730..fc07bbc2 100644 ---- a/third_party/stablehlo/temporary.patch -+++ b/third_party/stablehlo/temporary.patch -@@ -1,17 +1,3 @@ --# Copyright 2026 The OpenXLA Authors. --# --# Licensed under the Apache License, Version 2.0 (the "License"); --# you may not use this file except in compliance with the License. --# You may obtain a copy of the License at --# --# http://www.apache.org/licenses/LICENSE-2.0 --# --# Unless required by applicable law or agreed to in writing, software --# distributed under the License is distributed on an "AS IS" BASIS, --# WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied. --# See the License for the specific language governing permissions and --# limitations under the License. --# ============================================================================== - diff --ruN a/stablehlo/stablehlo/dialect/Serialization.cpp b/stablehlo/stablehlo/dialect/Serialization.cpp - --- stablehlo/stablehlo/dialect/Serialization.cpp - +++ stablehlo/stablehlo/dialect/Serialization.cpp diff --git a/third_party/xla/third_party/shardy/workspace.bzl b/third_party/xla/third_party/shardy/workspace.bzl index 89717bf9ab33df..cf3a5abc14b959 100644 --- a/third_party/xla/third_party/shardy/workspace.bzl +++ b/third_party/xla/third_party/shardy/workspace.bzl @@ -18,8 +18,8 @@ load("//third_party:repo.bzl", "tf_http_archive", "tf_mirror_urls") def repo(): - SHARDY_COMMIT = "eb23a98329aa70d991aa2d8a51a209af1f8df8fc" - SHARDY_SHA256 = "7cfbfda2f36c917d471a59cfed7ec9776adb96b6c2de3a0d7fd37f21dade36cf" + SHARDY_COMMIT = "0e66b80fc4aebf260f9b19a5fd9a71c5850f985f" + SHARDY_SHA256 = "a11e95eada3bd2daf1c1a0f1c6658f4ec4e27680ab58fbdf31eb24fa2ef9530d" tf_http_archive( name = "shardy", diff --git a/third_party/xla/third_party/triton/common/llvm_cl974500093.patch b/third_party/xla/third_party/triton/common/llvm_cl974500093.patch new file mode 100644 index 00000000000000..fbf632397ba433 --- /dev/null +++ b/third_party/xla/third_party/triton/common/llvm_cl974500093.patch @@ -0,0 +1,198 @@ +--- a/test/Conversion/amd/async-ops-alias-scopes.mlir ++++ b/test/Conversion/amd/async-ops-alias-scopes.mlir +@@ -16,7 +16,7 @@ + %ptr = tt.splat %arg0 : !tt.ptr -> tensor<64x1x!tt.ptr, #blocked> + %mask = tt.splat %maskVal : i1 -> tensor<64x1xi1, #blocked> + +- // COMMON: rocdl.global.load.async.lds {{.*}} {alias_scopes = [[[$ASYNC_COPY_SCOPE]]] ++ // COMMON: rocdl.global.load.async.lds {{.*}} <{alias_scopes = [[[$ASYNC_COPY_SCOPE]]] + // Check that store for 'other' has alias information set + // COMMON: llvm.store {{.*}} {alias_scopes = [[[$LOCAL_LOAD_SCOPE]]], {{.*}}, noalias_scopes = [[[$ASYNC_COPY_SCOPE]]] + %0 = ttg.async_copy_global_to_local %ptr, %arg1 mask %mask other %other : tensor<64x1x!tt.ptr, #blocked> -> <64x1xf32, #shared, #smem, mutable> +@@ -29,6 +29,7 @@ + // ----- + + // COMMON: [[$ASYNC_COPY_SCOPE:#.*]] = #llvm.alias_scope + #shared = #ttg.swizzled_shared<{vec = 1, perPhase = 1, maxPhase = 1, order = [1, 0]}> + #smem = #ttg.shared_memory +@@ -41,7 +42,7 @@ + %mask = tt.splat %maskVal : i1 -> tensor<8x64xi1, #blocked> + %other = arith.constant dense<1.000000e+00> : tensor<8x64xf32, #blocked> + +- // COMMON: rocdl.raw.ptr.buffer.load.async.lds {{.*}} {alias_scopes = [[[$ASYNC_COPY_SCOPE]]] ++ // COMMON: rocdl.raw.ptr.buffer.load.async.lds {{.*}} <{alias_scopes = [[[$ASYNC_COPY_SCOPE]]] + // Check that store for 'other' has alias information set + // COMMON: llvm.store {{.*}} {alias_scopes = [[[$LOCAL_LOAD_SCOPE]]], {{.*}}, noalias_scopes = [[[$ASYNC_COPY_SCOPE]]] + %65 = amdg.buffer_load_to_local %arg1[%arg2] mask=%mask other=%other into %arg3 : [tensor<8x64xi32, #blocked>] tensor<8x64xf32, #blocked> -> <8x64xf32, #shared, #smem, mutable> +@@ -75,7 +76,7 @@ + + // Test lowering path in AMD's MemoryOpToLLVM pattern + // GFX942: llvm.load {{.*}} {alias_scopes = [[[$LOCAL_LOAD_SCOPE]]], noalias_scopes = [[[$ASYNC_COPY_SCOPE]]] +- // GFX950: rocdl.ds.read.tr16.b64 {{.*}} {alias_scopes = [[[$LOCAL_LOAD_SCOPE]]], noalias_scopes = [[[$ASYNC_COPY_SCOPE]]] ++ // GFX950: rocdl.ds.read.tr16.b64 {{.*}} <{alias_scopes = [[[$LOCAL_LOAD_SCOPE]]], noalias_scopes = [[[$ASYNC_COPY_SCOPE]]] + %5 = ttg.local_load %arg2 token %3 : !ttg.memdesc<16x16xf16, #shared, #smem, mutable> -> tensor<16x16xf16, #ttg.dot_op<{opIdx = 0, parent = #mma, kWidth = 8}>> + + // Stores to keep the local_loads +@@ -107,7 +108,7 @@ + // We need the splat to allow the AxisAnalysis to work during lowering + %ptr = tt.splat %arg0 : !tt.ptr -> tensor<64x1x!tt.ptr, #blocked> + +- // COMMON: rocdl.global.load.async.lds {{.*}} {alias_scopes = [[[$ASYNC_COPY_SCOPE]]] ++ // COMMON: rocdl.global.load.async.lds {{.*}} <{alias_scopes = [[[$ASYNC_COPY_SCOPE]]] + %0 = ttg.async_copy_global_to_local %ptr, %arg1 : tensor<64x1x!tt.ptr, #blocked> -> <64x1xf32, #shared, #smem, mutable> + %1 = ttg.async_commit_group tokens %0 + + +--- a/test/Conversion/tritongpu_to_llvm.mlir ++++ b/test/Conversion/tritongpu_to_llvm.mlir +@@ -929,8 +929,8 @@ + tt.func @convert_dot_ldmatrix(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { + %AA = ttg.local_alloc %A : (tensor<16x16xf16, #blocked0>) -> !ttg.memdesc<16x16xf16, #shared0, #smem> + %BB = ttg.local_alloc %B : (tensor<16x16xf16, #blocked0>) -> !ttg.memdesc<16x16xf16, #shared0, #smem> +- // CHECK: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, num = 4 : i32, shape = #nvvm.ld_st_matrix_shape} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> +- // CHECK: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, num = 4 : i32, shape = #nvvm.ld_st_matrix_shape} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> ++ // CHECK: nvvm.ldmatrix %{{.*}}, num = 4, layout = , shape = , element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> ++ // CHECK: nvvm.ldmatrix %{{.*}}, num = 4, layout = , shape = , element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> + // CHECK-NOT: nvvm.ldmatrix + %AA_DOT = ttg.local_load %AA : !ttg.memdesc<16x16xf16, #shared0, #smem> -> tensor<16x16xf16, #dot_operand_a> + %BB_DOT = ttg.local_load %BB : !ttg.memdesc<16x16xf16, #shared0, #smem> -> tensor<16x16xf16, #dot_operand_b> +@@ -959,8 +959,8 @@ + tt.func @convert_dot_ldmatrix_swizzle(%A: tensor<16x16xf16, #blocked0>, %B: tensor<16x16xf16, #blocked0>) { + %AA = ttg.local_alloc %A : (tensor<16x16xf16, #blocked0>) -> !ttg.memdesc<16x16xf16, #shared0, #smem> + %BB = ttg.local_alloc %B : (tensor<16x16xf16, #blocked0>) -> !ttg.memdesc<16x16xf16, #shared0, #smem> +- // CHECK: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, num = 4 : i32, shape = #nvvm.ld_st_matrix_shape} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> +- // CHECK: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, num = 4 : i32, shape = #nvvm.ld_st_matrix_shape} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> ++ // CHECK: nvvm.ldmatrix %{{.*}}, num = 4, layout = , shape = , element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> ++ // CHECK: nvvm.ldmatrix %{{.*}}, num = 4, layout = , shape = , element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> + // CHECK-NOT: nvvm.ldmatrix + %AA_DOT = ttg.local_load %AA : !ttg.memdesc<16x16xf16, #shared0, #smem> -> tensor<16x16xf16, #dot_operand_a> + %BB_DOT = ttg.local_load %BB : !ttg.memdesc<16x16xf16, #shared0, #smem> -> tensor<16x16xf16, #dot_operand_b> +@@ -1041,7 +1041,7 @@ + tt.func @convert_dot_fp8(%A: tensor<16x16xf8E5M2, #blocked0>, %B: tensor<16x16xf8E5M2, #blocked0>) { + %AA = ttg.local_alloc %A : (tensor<16x16xf8E5M2, #blocked0>) -> !ttg.memdesc<16x16xf8E5M2, #shared0, #smem> + %BB = ttg.local_alloc %B : (tensor<16x16xf8E5M2, #blocked0>) -> !ttg.memdesc<16x16xf8E5M2, #shared0, #smem> +- // CHECK: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, num = 2 : i32, shape = #nvvm.ld_st_matrix_shape} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)> ++ // CHECK: nvvm.ldmatrix %{{.*}}, num = 2, layout = , shape = , element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32)> + // CHECK-NOT: nvvm.ldmatrix + %AA_DOT = ttg.local_load %AA : !ttg.memdesc<16x16xf8E5M2, #shared0, #smem> -> tensor<16x16xf8E5M2, #dot_operand_a> + %BB_DOT = ttg.local_load %BB : !ttg.memdesc<16x16xf8E5M2, #shared0, #smem> -> tensor<16x16xf8E5M2, #dot_operand_b> + +--- a/test/Conversion/tritongpu_to_llvm_blackwell.mlir ++++ b/test/Conversion/tritongpu_to_llvm_blackwell.mlir +@@ -963,7 +963,7 @@ + + // CHECK-LABEL: max_reduction + // CHECK: %[[M:.+]] = llvm.mlir.constant(-1 : i32) : i32 +-// CHECK: nvvm.redux.sync fmax %{{.*}}, %[[M]] {nan = true} : f32 -> f32 ++// CHECK: nvvm.redux.sync fmax %{{.*}}, %[[M]] nan = true : f32 -> f32 + // CHECK: nvvm.barrier + // CHECK: nvvm.shfl.sync bfly + // CHECK: nvvm.shfl.sync bfly +@@ -1008,7 +1008,7 @@ + // CHECK-LABEL: lower_ldmatrix_trans_b8 + tt.func @lower_ldmatrix_trans_b8(%A: !ttg.memdesc<128x64xf8E4M3FN, #shared, #smem, mutable, 1x128x64>) { + %0 = ttg.local_load %A : !ttg.memdesc<128x64xf8E4M3FN, #shared, #smem, mutable, 1x128x64> -> tensor<128x64xf8E4M3FN, #ttg.dot_op<{opIdx = 1, parent = #mma, kWidth = 4}>> +- // CHECK-COUNT-16: nvvm.ldmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout{{.*}}} : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> ++ // CHECK-COUNT-16: nvvm.ldmatrix %{{.*}}layout = {{.*}}element_type = : (!llvm.ptr<3>) -> !llvm.struct<(i32, i32, i32, i32)> + tt.return + } + } +@@ -1021,7 +1021,7 @@ + module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + // CHECK-LABEL: @stmatrix_b8_trans_linear + tt.func public @stmatrix_b8_trans_linear(%data: tensor<1x1x1x16x256xf8E4M3FN, #linear3>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout{{.*}}} : !llvm.ptr<3>, i32, i32, i32, i32 ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}layout = {{.*}}element_type = : !llvm.ptr<3>, i32, i32, i32, i32 + %0 = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<1x1x1x16x256xf8E4M3FN, #shared, #smem, mutable> + ttg.local_store %data, %0 : tensor<1x1x1x16x256xf8E4M3FN, #linear3> -> !ttg.memdesc<1x1x1x16x256xf8E4M3FN, #shared, #smem, mutable> + tt.return + +--- a/test/Conversion/tritongpu_to_llvm_hopper.mlir ++++ b/test/Conversion/tritongpu_to_llvm_hopper.mlir +@@ -441,7 +441,7 @@ + // CHECK-LABEL: distribute_to_shared_st_matrix_local_store + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @distribute_to_shared_st_matrix_local_store(%a: tensor<64x128xf16, #linear>) { +- // CHECK-COUNT-8: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-8: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<64x128xf16, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<64x128xf16, #linear> -> !ttg.memdesc<64x128xf16, #shared, #smem, mutable> +@@ -473,7 +473,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_local_store + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_local_store(%a: tensor<64x32xf16, #linear>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<64x32xf16, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<64x32xf16, #linear> -> !ttg.memdesc<64x32xf16, #shared, #smem, mutable> +@@ -495,7 +495,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_local_store + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_local_store(%a: tensor<32x32xf16, #linear>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<32x32xf16, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<32x32xf16, #linear> -> !ttg.memdesc<32x32xf16, #shared, #smem, mutable> +@@ -511,7 +511,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_x2_local_store_fp8 + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_x2_local_store_fp8(%a: tensor<64x16xf8E4M3FNUZ, #linear>) { +- // CHECK-COUNT-1: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-1: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<64x16xf8E4M3FNUZ, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<64x16xf8E4M3FNUZ, #linear> -> !ttg.memdesc<64x16xf8E4M3FNUZ, #shared, #smem, mutable> +@@ -527,7 +527,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_local_store_fp32 + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_local_store_fp32(%a: tensor<64x16xf32, #linear>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<64x16xf32, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<64x16xf32, #linear> -> !ttg.memdesc<64x16xf32, #shared, #smem, mutable> +@@ -544,7 +544,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_trans_local_store + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_trans_local_store(%a: tensor<64x32xf16, #linear>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<64x32xf16, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<64x32xf16, #linear> -> !ttg.memdesc<64x32xf16, #shared, #smem, mutable> +@@ -566,7 +566,7 @@ + // CHECK-LABEL: linear_to_swizzled_st_matrix_trans_local_store + module attributes {"ttg.target" = "cuda:90", "ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32, "ttg.threads-per-warp" = 32 : i32} { + tt.func @linear_to_swizzled_st_matrix_trans_local_store(%a: tensor<16x32xf16, #linear>) { +- // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} {eltType = #nvvm.ld_st_matrix_elt_type, layout = #nvvm.mma_layout, shape = #nvvm.ld_st_matrix_shape} ++ // CHECK-COUNT-2: nvvm.stmatrix %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}}, %{{.*}} layout = , shape = , element_type = + // CHECK: llvm.return + %b = ttg.local_alloc {allocation.offset = 0 : i32} : () -> !ttg.memdesc<16x32xf16, #shared, #smem, mutable> + ttg.local_store %a, %b : tensor<16x32xf16, #linear> -> !ttg.memdesc<16x32xf16, #shared, #smem, mutable> + +--- a/test/Conversion/tritonnvidiagpu_to_llvm.mlir ++++ b/test/Conversion/tritonnvidiagpu_to_llvm.mlir +@@ -514,7 +514,7 @@ + + module attributes {"ttg.num-ctas" = 1 : i32, "ttg.num-warps" = 4 : i32} { + // CHECK-LABEL: async_tma_store_wait_read_only +- // CHECK: nvvm.cp.async.bulk.wait_group 0 {read} ++ // CHECK: nvvm.cp.async.bulk.wait_group 0 read + tt.func @async_tma_store_wait_read_only() { + ttng.async_tma_store_wait {pendings = 0 : i32, read_only} + tt.return +@@ -650,7 +650,7 @@ + tt.func public @async_copy_mbarrier_arrive(%arg0: !ttg.memdesc<1xi64, #shared, #ttg.shared_memory>) attributes { noinline = false } { + // CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} : !llvm.ptr<3> + ttng.async_copy_mbarrier_arrive %arg0 : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory> +- // CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} {noinc = true} : !llvm.ptr<3> ++ // CHECK: nvvm.cp.async.mbarrier.arrive %{{.*}} noinc = true : !llvm.ptr<3> + ttng.async_copy_mbarrier_arrive %arg0 { noIncrement } : !ttg.memdesc<1xi64, #shared, #ttg.shared_memory> + tt.return + } + diff --git a/third_party/xla/third_party/triton/common/series.bzl b/third_party/xla/third_party/triton/common/series.bzl index f8bb0562969056..bf20e6d451f321 100644 --- a/third_party/xla/third_party/triton/common/series.bzl +++ b/third_party/xla/third_party/triton/common/series.bzl @@ -55,5 +55,6 @@ common_patch_list = [ "//third_party/triton:common/blackwell_nvfp4_mn_major_fallback.patch", "//third_party/triton:common/llvm_cl966791199.patch", "//third_party/triton:common/llvm_cl969980821.patch", + "//third_party/triton:common/llvm_cl974500093.patch", # Add new patches just above this line ] diff --git a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/convert_s8_s32.hlo b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/convert_s8_s32.hlo index 68df2c04e7acb7..5709a6d143fab6 100644 --- a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/convert_s8_s32.hlo +++ b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/convert_s8_s32.hlo @@ -42,7 +42,7 @@ fusion { // CHECK-VECTOR-LABEL: func.func @wrapped_fusion_impl // CHECK-VECTOR: %[[ALLOCA_0:.*]] = memref.alloca() : memref<1x4xi8> // CHECK-VECTOR: %[[ALLOCA_1:.*]] = memref.alloca() : memref<1x4xi8> -// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() {alignment = 64 : i64} : memref<1x4xi32> +// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() alignment = 64 : memref<1x4xi32> // CHECK-VECTOR: memref.copy %{{.*}}, %[[ALLOCA_0]] // NEW-VECTOR-LABEL: func.func @wrapped_fusion_impl diff --git a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/is_finite.hlo b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/is_finite.hlo index a06cc072c12087..cdb8e7604f65f0 100644 --- a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/is_finite.hlo +++ b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/is_finite.hlo @@ -31,7 +31,7 @@ // CHECK-VECTOR-SAME: %[[OUT:[^:]+]]: memref<1x9xi8>, // CHECK-VECTOR-SAME: %[[PID:[^:]+]]: index) // CHECK-VECTOR: %[[ALLOCA_ARG:.*]] = memref.alloca() : memref<1x16xf16> -// CHECK-VECTOR: %[[ALLOCA_OUT:.*]] = memref.alloca() {{{.*}}} : memref<1x16xi8> +// CHECK-VECTOR: %[[ALLOCA_OUT:.*]] = memref.alloca() alignment = 64 : memref<1x16xi8> // CHECK-VECTOR: memref.copy %[[ARG]], // CHECK-VECTOR: scf.for // CHECK-VECTOR: %[[READ:.*]] = vector.transfer_read %[[ALLOCA_ARG]]{{.*}} : memref<1x16xf16>, vector<8xf16> diff --git a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/reduce_precision.hlo b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/reduce_precision.hlo index ec1dad0f5e36ee..b59d9502b3b75d 100644 --- a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/reduce_precision.hlo +++ b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/elementwise/reduce_precision.hlo @@ -45,7 +45,7 @@ fusion { // CHECK-VECTOR-LABEL: func.func @wrapped_fusion_impl // CHECK-VECTOR: %[[ALLOCA_0:.*]] = memref.alloca() : memref<1x4xf32> // CHECK-VECTOR: %[[ALLOCA_1:.*]] = memref.alloca() : memref<1x4xf32> -// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() {alignment = 64 : i64} : memref<1x4xf32> +// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() alignment = 64 : memref<1x4xf32> // CHECK-VECTOR: memref.copy %{{.*}}, %[[ALLOCA_0]] // CHECK-VECTOR: memref.copy %{{.*}}, %{{.*}} diff --git a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/transpose/transpose_c64.hlo b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/transpose/transpose_c64.hlo index 2449687fc3fe64..8d5a29f7f8c1fb 100644 --- a/third_party/xla/xla/backends/cpu/codegen/tiled/tests/transpose/transpose_c64.hlo +++ b/third_party/xla/xla/backends/cpu/codegen/tiled/tests/transpose/transpose_c64.hlo @@ -41,7 +41,7 @@ fusion { // CHECK-VECTOR-LABEL: func.func @wrapped_fusion_impl // CHECK-VECTOR: %[[ALLOCA_0:.*]] = memref.alloca() : memref<16x1x1x16xcomplex> // CHECK-VECTOR: %[[ALLOCA_1:.*]] = memref.alloca() : memref<16x1x1x16xcomplex> -// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() {alignment = 64 : i64} : memref<16x1x1x16xcomplex> +// CHECK-VECTOR: %[[ALLOCA_2:.*]] = memref.alloca() alignment = 64 : memref<16x1x1x16xcomplex> // CHECK-VECTOR: memref.copy %{{.*}}, %[[ALLOCA_0]] // CHECK-VECTOR: memref.copy %[[ALLOCA_2]], %{{.*}} diff --git a/third_party/xla/xla/backends/cpu/codegen/tiled/transforms/tests/vectorize_xtile.mlir b/third_party/xla/xla/backends/cpu/codegen/tiled/transforms/tests/vectorize_xtile.mlir index afdf033e7b769a..bf921c30019fcb 100644 --- a/third_party/xla/xla/backends/cpu/codegen/tiled/transforms/tests/vectorize_xtile.mlir +++ b/third_party/xla/xla/backends/cpu/codegen/tiled/transforms/tests/vectorize_xtile.mlir @@ -967,7 +967,7 @@ func.func @test_slice(%arg0: tensor<8x16xf32>) -> tensor<4x8xf32> { } // CHECK-LABEL: @test_slice // CHECK: %[[CAST:.*]] = builtin.unrealized_conversion_cast %{{.*}} : tensor<8x16xf32> to vector<8x16xf32> -// CHECK: %[[SLICE:.*]] = vector.extract_strided_slice %[[CAST]] {offsets = [2, 4], sizes = [4, 8], strides = [1, 1]} : vector<8x16xf32> to vector<4x8xf32> +// CHECK: %[[SLICE:.*]] = vector.extract_strided_slice %[[CAST]] offsets = [2, 4], sizes = [4, 8], strides = [1, 1] : vector<8x16xf32> to vector<4x8xf32> // CHECK: %[[RET:.*]] = builtin.unrealized_conversion_cast %[[SLICE]] : vector<4x8xf32> to tensor<4x8xf32> // CHECK: return %[[RET]] @@ -998,8 +998,8 @@ func.func @test_concatenate(%arg0: tensor<4x8xf32>, %arg1: tensor<4x8xf32>) -> t // CHECK-DAG: %[[CST:.*]] = arith.constant dense<0.000000e+00> : vector<4x16xf32> // CHECK-DAG: %[[LHS:.*]] = builtin.unrealized_conversion_cast %[[ARG0]] : tensor<4x8xf32> to vector<4x8xf32> // CHECK-DAG: %[[RHS:.*]] = builtin.unrealized_conversion_cast %[[ARG1]] : tensor<4x8xf32> to vector<4x8xf32> -// CHECK: %[[INSERT0:.*]] = vector.insert_strided_slice %[[LHS]], %[[CST]] {offsets = [0, 0], strides = [1, 1]} : vector<4x8xf32> into vector<4x16xf32> -// CHECK: %[[INSERT1:.*]] = vector.insert_strided_slice %[[RHS]], %[[INSERT0]] {offsets = [0, 8], strides = [1, 1]} : vector<4x8xf32> into vector<4x16xf32> +// CHECK: %[[INSERT0:.*]] = vector.insert_strided_slice %[[LHS]], %[[CST]] offsets = [0, 0], strides = [1, 1] : vector<4x8xf32> into vector<4x16xf32> +// CHECK: %[[INSERT1:.*]] = vector.insert_strided_slice %[[RHS]], %[[INSERT0]] offsets = [0, 8], strides = [1, 1] : vector<4x8xf32> into vector<4x16xf32> // CHECK: %[[RET:.*]] = builtin.unrealized_conversion_cast %[[INSERT1]] : vector<4x16xf32> to tensor<4x16xf32> // CHECK: return %[[RET]] diff --git a/third_party/xla/xla/backends/gpu/codegen/emitters/transforms/tests/promote_shuffle_to_dpp.mlir b/third_party/xla/xla/backends/gpu/codegen/emitters/transforms/tests/promote_shuffle_to_dpp.mlir index 3a8585135dd907..3967d2a6830159 100644 --- a/third_party/xla/xla/backends/gpu/codegen/emitters/transforms/tests/promote_shuffle_to_dpp.mlir +++ b/third_party/xla/xla/backends/gpu/codegen/emitters/transforms/tests/promote_shuffle_to_dpp.mlir @@ -25,7 +25,7 @@ module { // CHECK-LABEL: @shuffle_down_1 // CHECK-SAME: (%[[ARG:.*]]: f32) -// CHECK: %[[DPP:.*]] = amdgpu.dpp %[[ARG]] %[[ARG]] row_shl(1 : i32) {bound_ctrl = true} : f32 +// CHECK: %[[DPP:.*]] = amdgpu.dpp %[[ARG]] %[[ARG]] row_shl(1 : i32) bound_ctrl(true) : f32 // CHECK: return %[[DPP]] : f32 // ----- @@ -40,7 +40,7 @@ module { } // CHECK-LABEL: @shuffle_down_4 -// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(4 : i32) {bound_ctrl = true} : f32 +// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(4 : i32) bound_ctrl(true) : f32 // ----- @@ -54,7 +54,7 @@ module { } // CHECK-LABEL: @shuffle_down_8 -// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(8 : i32) {bound_ctrl = true} : f32 +// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(8 : i32) bound_ctrl(true) : f32 // ----- @@ -68,7 +68,7 @@ module { } // CHECK-LABEL: @shuffle_down_15 -// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(15 : i32) {bound_ctrl = true} : f32 +// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(15 : i32) bound_ctrl(true) : f32 // ----- @@ -134,7 +134,7 @@ module { } // CHECK-LABEL: @shuffle_down_i32 -// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(2 : i32) {bound_ctrl = true} : i32 +// CHECK: amdgpu.dpp %{{.*}} %{{.*}} row_shl(2 : i32) bound_ctrl(true) : i32 // ----- diff --git a/third_party/xla/xla/codegen/emitters/transforms/tests/lower_xla_intrinsic_lib.mlir b/third_party/xla/xla/codegen/emitters/transforms/tests/lower_xla_intrinsic_lib.mlir index c66b8fd4da6f61..1f697977a957c0 100644 --- a/third_party/xla/xla/codegen/emitters/transforms/tests/lower_xla_intrinsic_lib.mlir +++ b/third_party/xla/xla/codegen/emitters/transforms/tests/lower_xla_intrinsic_lib.mlir @@ -279,18 +279,18 @@ module { // CHECK-LABEL: @tanh_f32_vector_64 // CHECK-NOT: math.tanh // CHECK: %[[INIT:.*]] = arith.constant dense<0.000000e+00> : vector<64xf32> -// CHECK: %[[S0:.*]] = vector.extract_strided_slice %arg0 {offsets = [0], sizes = [16], strides = [1]} : vector<64xf32> to vector<16xf32> +// CHECK: %[[S0:.*]] = vector.extract_strided_slice %arg0 offsets = [0], sizes = [16], strides = [1] : vector<64xf32> to vector<16xf32> // CHECK: %[[C0:.*]] = call @xla.tanh.v16f32(%[[S0]]) -// CHECK: %[[I0:.*]] = vector.insert_strided_slice %[[C0]], %[[INIT]] {offsets = [0], strides = [1]} : vector<16xf32> into vector<64xf32> -// CHECK: %[[S1:.*]] = vector.extract_strided_slice %arg0 {offsets = [16], sizes = [16], strides = [1]} : vector<64xf32> to vector<16xf32> +// CHECK: %[[I0:.*]] = vector.insert_strided_slice %[[C0]], %[[INIT]] offsets = [0], strides = [1] : vector<16xf32> into vector<64xf32> +// CHECK: %[[S1:.*]] = vector.extract_strided_slice %arg0 offsets = [16], sizes = [16], strides = [1] : vector<64xf32> to vector<16xf32> // CHECK: %[[C1:.*]] = call @xla.tanh.v16f32(%[[S1]]) -// CHECK: %[[I1:.*]] = vector.insert_strided_slice %[[C1]], %[[I0]] {offsets = [16], strides = [1]} : vector<16xf32> into vector<64xf32> -// CHECK: %[[S2:.*]] = vector.extract_strided_slice %arg0 {offsets = [32], sizes = [16], strides = [1]} : vector<64xf32> to vector<16xf32> +// CHECK: %[[I1:.*]] = vector.insert_strided_slice %[[C1]], %[[I0]] offsets = [16], strides = [1] : vector<16xf32> into vector<64xf32> +// CHECK: %[[S2:.*]] = vector.extract_strided_slice %arg0 offsets = [32], sizes = [16], strides = [1] : vector<64xf32> to vector<16xf32> // CHECK: %[[C2:.*]] = call @xla.tanh.v16f32(%[[S2]]) -// CHECK: %[[I2:.*]] = vector.insert_strided_slice %[[C2]], %[[I1]] {offsets = [32], strides = [1]} : vector<16xf32> into vector<64xf32> -// CHECK: %[[S3:.*]] = vector.extract_strided_slice %arg0 {offsets = [48], sizes = [16], strides = [1]} : vector<64xf32> to vector<16xf32> +// CHECK: %[[I2:.*]] = vector.insert_strided_slice %[[C2]], %[[I1]] offsets = [32], strides = [1] : vector<16xf32> into vector<64xf32> +// CHECK: %[[S3:.*]] = vector.extract_strided_slice %arg0 offsets = [48], sizes = [16], strides = [1] : vector<64xf32> to vector<16xf32> // CHECK: %[[C3:.*]] = call @xla.tanh.v16f32(%[[S3]]) -// CHECK: %[[I3:.*]] = vector.insert_strided_slice %[[C3]], %[[I2]] {offsets = [48], strides = [1]} : vector<16xf32> into vector<64xf32> +// CHECK: %[[I3:.*]] = vector.insert_strided_slice %[[C3]], %[[I2]] offsets = [48], strides = [1] : vector<16xf32> into vector<64xf32> // CHECK: return %[[I3]] : vector<64xf32> // ----- @@ -304,11 +304,11 @@ module { // CHECK-LABEL: @tanh_bf16_vector_32 // CHECK-NOT: math.tanh -// CHECK: %[[S0:.*]] = vector.extract_strided_slice %arg0 {offsets = [0], sizes = [16], strides = [1]} : vector<32xbf16> to vector<16xbf16> +// CHECK: %[[S0:.*]] = vector.extract_strided_slice %arg0 offsets = [0], sizes = [16], strides = [1] : vector<32xbf16> to vector<16xbf16> // CHECK: %[[EXT0:.*]] = arith.extf %[[S0]] : vector<16xbf16> to vector<16xf32> // CHECK: %[[C0:.*]] = call @xla.tanh.v16f32(%[[EXT0]]) // CHECK: %[[TR0:.*]] = call @xla.fptrunc.v8f32.to.v8bf16 -// CHECK: %[[S1:.*]] = vector.extract_strided_slice %arg0 {offsets = [16], sizes = [16], strides = [1]} : vector<32xbf16> to vector<16xbf16> +// CHECK: %[[S1:.*]] = vector.extract_strided_slice %arg0 offsets = [16], sizes = [16], strides = [1] : vector<32xbf16> to vector<16xbf16> // CHECK: %[[EXT1:.*]] = arith.extf %[[S1]] : vector<16xbf16> to vector<16xf32> // CHECK: %[[C1:.*]] = call @xla.tanh.v16f32(%[[EXT1]]) // CHECK: %[[TR1:.*]] = call @xla.fptrunc.v8f32.to.v8bf16 diff --git a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/memref/dim.mlir b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/memref/dim.mlir index 26cae97bab729d..a26d6ea9b83c1f 100644 --- a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/memref/dim.mlir +++ b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/memref/dim.mlir @@ -15,7 +15,7 @@ // RUN: mlir-interpreter-runner %s -run-all | FileCheck %s func.func @dim() -> index { - %alloc = memref.alloc() {alignment = 64 : i64} : memref<10x50xf32> + %alloc = memref.alloc() alignment = 64 : memref<10x50xf32> %c1 = arith.constant 1 : index %dim = memref.dim %alloc, %c1 : memref<10x50xf32> return %dim : index diff --git a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/extract_strided_slice.mlir b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/extract_strided_slice.mlir index 68b934c24a0c69..af7b9265a36f41 100644 --- a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/extract_strided_slice.mlir +++ b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/extract_strided_slice.mlir @@ -18,12 +18,12 @@ func.func @extract_strided_slice() -> vector<2x3xi32> { %c = arith.constant dense<[[1,2,3,4], [5,6,7,8], [9,10,11,12]]> : vector<3x4xi32> - %o = vector.extract_strided_slice %c { + %o = vector.extract_strided_slice %c offsets = [0, 1], sizes = [2, 3], // TODO(jreiffers): Test non-unit strides when supported by verifier. strides = [1, 1] - } : vector<3x4xi32> to vector<2x3xi32> + : vector<3x4xi32> to vector<2x3xi32> return %o : vector<2x3xi32> } diff --git a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/insert_strided_slice.mlir b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/insert_strided_slice.mlir index 3f0df0acf7a9f8..d8034ce3655282 100644 --- a/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/insert_strided_slice.mlir +++ b/third_party/xla/xla/mlir/tools/mlir_interpreter/dialects/tests/vector/insert_strided_slice.mlir @@ -17,11 +17,11 @@ func.func @insert_strided_slice() -> (vector<3x4xi32>, vector<3x4xi32>) { %v = arith.constant dense<[[2, 3, 4], [6, 7, 8]]> : vector<2x3xi32> %c = arith.constant dense<0> : vector<3x4xi32> - %o = vector.insert_strided_slice %v, %c { + %o = vector.insert_strided_slice %v, %c offsets = [0, 1], // TODO(jreiffers): Test non-unit strides when supported by verifier. strides = [1, 1] - } : vector<2x3xi32> into vector<3x4xi32> + : vector<2x3xi32> into vector<3x4xi32> return %c, %o : vector<3x4xi32>, vector<3x4xi32> } diff --git a/third_party/xla/xla/mlir_hlo/tests/Dialect/deallocation/buffer_reuse.mlir b/third_party/xla/xla/mlir_hlo/tests/Dialect/deallocation/buffer_reuse.mlir index 08b263553fa5ce..683d7fa35913ab 100644 --- a/third_party/xla/xla/mlir_hlo/tests/Dialect/deallocation/buffer_reuse.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/Dialect/deallocation/buffer_reuse.mlir @@ -453,14 +453,14 @@ func.func @hoist_from_if(%cond: i1) { // ----- func.func @propagate_alignment_attr() { - %alloc = memref.alloc() {alignment = 64 : i64} : memref + %alloc = memref.alloc() alignment = 64 : memref "test.use"(%alloc) : (memref) -> () memref.dealloc %alloc : memref return } // CHECK-LABEL: @propagate_alignment_attr -// CHECK-NEXT: memref.alloca() {alignment = 64 : i64} : memref +// CHECK-NEXT: memref.alloca() alignment = 64 : memref // ----- diff --git a/third_party/xla/xla/mlir_hlo/tests/alloc_to_arg.mlir b/third_party/xla/xla/mlir_hlo/tests/alloc_to_arg.mlir index d8f809329fd2e1..f48508c80bb999 100644 --- a/third_party/xla/xla/mlir_hlo/tests/alloc_to_arg.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/alloc_to_arg.mlir @@ -36,7 +36,7 @@ func.func @not_alloc(%arg0: memref<8xf32>) -> memref<8xf32> { func.func @fusion() -> memref<4x4x8x32xf32> { // CHECK: %[[COLLAPSE_SHAPE:.*]] = memref.collapse_shape %[[ARG0]] {{\[\[}}0, 1, 2], [3{{\]\]}} // CHECK: "some.use"(%[[COLLAPSE_SHAPE]], %[[ARG0]]) - %alloc = memref.alloc() {alignment = 64 : i64} : memref<128x32xf32> + %alloc = memref.alloc() alignment = 64 : memref<128x32xf32> %expand_shape = memref.expand_shape %alloc [[0, 1, 2], [3]] output_shape [4, 4, 8, 32] : memref<128x32xf32> into memref<4x4x8x32xf32> "some.use"(%alloc, %expand_shape) : (memref<128x32xf32>, memref<4x4x8x32xf32>) -> () return %expand_shape : memref<4x4x8x32xf32> diff --git a/third_party/xla/xla/mlir_hlo/tests/collapse_parallel_loops_to_1d_pass.mlir b/third_party/xla/xla/mlir_hlo/tests/collapse_parallel_loops_to_1d_pass.mlir index c1dfbde36b8155..90862c3afee412 100644 --- a/third_party/xla/xla/mlir_hlo/tests/collapse_parallel_loops_to_1d_pass.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/collapse_parallel_loops_to_1d_pass.mlir @@ -20,7 +20,7 @@ func.func @parallel_2d(%arg0: memref<4x4xf32>, %arg1: memref<4x4xf32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c4 = arith.constant 4 : index - %0 = memref.alloc() {alignment = 128 : i64} : memref<4x4xf32> + %0 = memref.alloc() alignment = 128 : memref<4x4xf32> scf.parallel (%arg2, %arg3) = (%c0, %c0) to (%c4, %c4) step (%c1, %c1) { // CHECK: scf.parallel ({{[^.]+}}) %2 = memref.load %arg0[%arg2,%arg3] : memref<4x4xf32> diff --git a/third_party/xla/xla/mlir_hlo/tests/naive_copy_removal.mlir b/third_party/xla/xla/mlir_hlo/tests/naive_copy_removal.mlir index d95a23f6b2330c..f269f8eb54aef8 100644 --- a/third_party/xla/xla/mlir_hlo/tests/naive_copy_removal.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/naive_copy_removal.mlir @@ -17,7 +17,7 @@ func.func @target_is_alloc(%arg0: memref<8x8xf32>) -> memref<8x8xf32> { %c4 = arith.constant 4 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> memref.copy %arg0, %alloc_4: memref<8x8xf32> to memref<8x8xf32> return %arg0 : memref<8x8xf32> } @@ -34,7 +34,7 @@ func.func @target_is_alloc_with_other_stores(%arg0: memref<8x8xf32>) -> memref<8x8xf32> { %c4 = arith.constant 4 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> memref.copy %arg0, %alloc_4: memref<8x8xf32> to memref<8x8xf32> linalg.fill ins(%cst_0 : f32) outs(%alloc_4 : memref<8x8xf32>) memref.store %cst_0, %alloc_4[%c4, %c4] : memref<8x8xf32> @@ -55,7 +55,7 @@ func.func @target_is_alloc_with_other_stores(%arg0: memref<8x8xf32>) func.func @target_is_subview(%arg0: memref<8x8xf32>) -> memref<8x8xf32> { %c4 = arith.constant 4 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> %subview_5 = memref.subview %alloc_4[0, 0] [%c4, %c4] [1, 1] : memref<8x8xf32> to memref> memref.copy %arg0, %subview_5 : @@ -75,7 +75,7 @@ func.func @target_is_subview_of_subview(%arg0: memref<8x8xf32>) -> memref<8x8xf32> { %c4 = arith.constant 4 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> %subview_5 = memref.subview %alloc_4[0, 0] [%c4, %c4] [1, 1] : memref<8x8xf32> to memref> %subview_6 = memref.subview %subview_5[0, 0] [%c4, %c4] [1, 1] : @@ -97,7 +97,7 @@ func.func @do_not_simplify_subview(%arg0: memref<8x8xf32>) -> vector<8x8xf32> { %c4 = arith.constant 4 : index %c0 = arith.constant 0 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> %subview_5 = memref.subview %alloc_4[0, 0] [%c4, %c4] [1, 1] : memref<8x8xf32> to memref> memref.copy %arg0, %subview_5 : @@ -119,7 +119,7 @@ func.func @do_not_simplify_alloc(%arg0: memref<8x8xf32>) -> vector<8x8xf32> { %c4 = arith.constant 4 : index %c0 = arith.constant 0 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> memref.copy %arg0, %alloc_4 : memref<8x8xf32> to memref<8x8xf32> %27 = vector.transfer_read %alloc_4[%c0, %c0], %cst_0 : memref<8x8xf32>, vector<8x8xf32> @@ -137,7 +137,7 @@ func.func @do_not_simplify_subview_with_other_use(%arg0: memref<8x8xf32>) -> memref<8x8xf32> { %c4 = arith.constant 4 : index %cst_0 = arith.constant 0.000000e+00 : f32 - %alloc_4 = memref.alloc() {alignment = 64 : i64} : memref<8x8xf32> + %alloc_4 = memref.alloc() alignment = 64 : memref<8x8xf32> %subview_5 = memref.subview %alloc_4[0, 0] [%c4, %c4] [1, 1] : memref<8x8xf32> to memref> %subview_6 = memref.subview %alloc_4[0, 0] [%c4, %c4] [1, 1] : diff --git a/third_party/xla/xla/mlir_hlo/tests/tile_loops.mlir b/third_party/xla/xla/mlir_hlo/tests/tile_loops.mlir index 8540d90ee6a76f..a8201cdc6e0180 100644 --- a/third_party/xla/xla/mlir_hlo/tests/tile_loops.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/tile_loops.mlir @@ -20,7 +20,7 @@ func.func @parallel_loop(%arg0: memref<16xf32>, %arg1: memref<16xf32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c16 = arith.constant 16 : index - %0 = memref.alloc() {alignment = 128 : i64} : memref<16xf32> + %0 = memref.alloc() alignment = 128 : memref<16xf32> scf.parallel (%arg2) = (%c0) to (%c16) step (%c1) { // CHECK-DAG: %[[C8:.*]] = arith.constant 8 // CHECK-DAG: %[[C4:.*]] = arith.constant 4 @@ -102,7 +102,7 @@ func.func @complex_access(%arg0: memref<16xf32>, %arg1: memref<4xf32>) { %c0 = arith.constant 0 : index %c1 = arith.constant 1 : index %c4 = arith.constant 4 : index - %0 = memref.alloc() {alignment = 128 : i64} : memref<4xf32> + %0 = memref.alloc() alignment = 128 : memref<4xf32> scf.parallel (%arg2) = (%c0) to (%c4) step (%c1) { // CHECK-DAG: %[[C2:.*]] = arith.constant 2 // CHECK: scf.parallel {{.*}} step (%[[C2]]) diff --git a/third_party/xla/xla/mlir_hlo/tests/vectorize_copy.mlir b/third_party/xla/xla/mlir_hlo/tests/vectorize_copy.mlir index cfb5df2f8eb8ce..a546e46f4604fa 100644 --- a/third_party/xla/xla/mlir_hlo/tests/vectorize_copy.mlir +++ b/third_party/xla/xla/mlir_hlo/tests/vectorize_copy.mlir @@ -44,7 +44,7 @@ func.func @do_not_vectorize_continuous_copy(%arg: memref<10x10xf32>) -> memref<1 func.func @tile_to_continuous_memref(%arg: memref<3x512xf32, strided<[768, 1]>>) -> (memref<3x512xf32>) { - %alloc = memref.alloc() {alignment = 64 : i64} : memref<3x512xf32> + %alloc = memref.alloc() alignment = 64 : memref<3x512xf32> memref.copy %arg, %alloc : memref<3x512xf32, strided<[768, 1]>> to memref<3x512xf32> return %alloc: memref<3x512xf32> } From 24585459c5bae7e5edd818afbfd7aa4775487697 Mon Sep 17 00:00:00 2001 From: Emily Fertig Date: Tue, 1 Sep 2026 14:50:54 -0700 Subject: [PATCH 41/43] Prefactor for supporting remote tile SPMEM transfers. PiperOrigin-RevId: 974720484 --- .../xla/xla/mosaic/dialect/tpu/tpu_dialect.cc | 37 ++++++++++--------- .../xla/xla/mosaic/dialect/tpu/tpu_dialect.h | 9 +++++ 2 files changed, 29 insertions(+), 17 deletions(-) diff --git a/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.cc b/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.cc index 3064f345c4710b..0b18d31190c21f 100644 --- a/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.cc +++ b/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.cc @@ -405,7 +405,6 @@ int64_t TiledLayoutAttr::getUntiledRank() const { return mlir::tpu::getUntiledRank(getTiles(), getRank()); } -namespace { FailureOr> getExpandedShape( const ArrayRef untiled_shape, const ArrayRef tiles, const bool require_alignment) { @@ -433,7 +432,6 @@ FailureOr> getExpandedShape( } return shape; } -} // namespace SmallVector TiledLayoutAttr::getContiguousTileStrides( const ArrayRef tiles, const ArrayRef shape) { @@ -510,27 +508,26 @@ SmallVector TiledLayoutAttr::getExpandedShape( /*require_alignment=*/false); } -SmallVector TiledLayoutAttr::getExpandedStrides() const { - if (getTiles().empty()) { - return SmallVector(getTileStrides()); +FailureOr> getExpandedStrides( + ArrayRef tiles, ArrayRef tile_strides) { + if (tiles.empty()) { + return SmallVector(tile_strides); } - SmallVector strides(getTileStrides()); + SmallVector strides(tile_strides); // Expand front tile - const xla::Tile& first_tile = getTiles().front(); - const FailureOr> failure_or_expanded_tile = - mlir::tpu::getExpandedShape(first_tile.dimensions(), - getTiles().drop_front(), - /*require_alignment=*/true); - // Verification should ensure this: - assert(succeeded(failure_or_expanded_tile)); - const SmallVector& expanded_tile = *failure_or_expanded_tile; - strides.resize_for_overwrite(getRank() + expanded_tile.size()); + const xla::Tile& first_tile = tiles.front(); + FAILUREOR_ASSIGN_OR_RETURN( + SmallVector expanded_tile, + mlir::tpu::getExpandedShape(first_tile.dimensions(), tiles.drop_front(), + /*require_alignment=*/true)); + const int64_t rank = tile_strides.size(); + strides.resize_for_overwrite(rank + expanded_tile.size()); int64_t first_tile_size = llvm::product_of(first_tile.dimensions()); int64_t tile_size = 1; for (int64_t d = strides.size() - 1; d >= 0; --d) { - if (d >= getRank()) { + if (d >= rank) { const int64_t new_stride = tile_size; - tile_size *= expanded_tile[d - getRank()]; + tile_size *= expanded_tile[d - rank]; strides[d] = new_stride; } else { if (ShapedType::isStatic(strides[d])) { @@ -541,6 +538,12 @@ SmallVector TiledLayoutAttr::getExpandedStrides() const { return strides; } +SmallVector TiledLayoutAttr::getExpandedStrides() const { + auto strides = mlir::tpu::getExpandedStrides(getTiles(), getTileStrides()); + CHECK(succeeded(strides)); + return *strides; +} + SmallVector TiledLayoutAttr::getSubtileUnit( const ArrayRef tiles) { assert(!tiles.empty()); diff --git a/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.h b/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.h index 3f886e4a1a2347..a87e9cd7426b5a 100644 --- a/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.h +++ b/third_party/xla/xla/mosaic/dialect/tpu/tpu_dialect.h @@ -95,6 +95,15 @@ LogicalResult specializeMemorySpace(TypedValue value, // vector ops. This functions inverts the layout erasure applied to the value. MemRefType getMemRefType(Value value); +// Expands a shape according to the tiles list, with optional alignment checks. +FailureOr> getExpandedShape( + ArrayRef untiled_shape, ArrayRef tiles, + bool require_alignment = false); + +// Evaluates the expanded strides for the given tiles and tile strides. +FailureOr> getExpandedStrides( + ArrayRef tiles, ArrayRef tile_strides); + // Returns the remainder of the given value when divided by the given divisor. // Returns nullopt if the remainder is not known. std::optional getRemainder(Value val, int64_t divisor, From e22f2b27decb17d5396811f71a7e1f45cdbc36c9 Mon Sep 17 00:00:00 2001 From: Bill Varcho Date: Tue, 1 Sep 2026 15:05:11 -0700 Subject: [PATCH 42/43] [SDY][Bug-Fix] Fix Shardy shard_map export for nested ManualComputationOp across calls PiperOrigin-RevId: 974728954 --- .../stablehlo_round_trip/shard_map_export.cc | 52 ++++++++++++------- ...stablehlo_round_trip_shard_map_export.mlir | 40 ++++++++++++++ 2 files changed, 72 insertions(+), 20 deletions(-) diff --git a/third_party/xla/xla/service/spmd/shardy/stablehlo_round_trip/shard_map_export.cc b/third_party/xla/xla/service/spmd/shardy/stablehlo_round_trip/shard_map_export.cc index bcec389dfcd402..3d2053b44eadc9 100644 --- a/third_party/xla/xla/service/spmd/shardy/stablehlo_round_trip/shard_map_export.cc +++ b/third_party/xla/xla/service/spmd/shardy/stablehlo_round_trip/shard_map_export.cc @@ -184,16 +184,25 @@ void setOpManualAxes(Operation* op, ManualAxesAttr manualAxes, op->setAttr(kManualAxes, manualAxes); } -void setFuncManualAxesRecursively(FuncOp funcOp, ManualAxesAttr manualAxes, - Attribute meshOrRef, - const mlir::SymbolTable& symbolTable); - -mlir::WalkResult setManualAxes(Operation* op, ManualAxesAttr manualAxes, - Attribute meshOrRef, - const mlir::SymbolTable& symbolTable) { - if (mlir::isa(op)) { - // Skip `ManualComputationOp`s and their nested operations, they will - // be handled separately. +void setManualAxesForOpsInBody( + ManualComputationOp op, const mlir::SymbolTable& symbolTable, + ManualComputationToParentManualAxes& parentManualCompAxes); + +void setFuncManualAxesRecursively( + FuncOp funcOp, ManualAxesAttr manualAxes, Attribute meshOrRef, + const mlir::SymbolTable& symbolTable, + ManualComputationToParentManualAxes& parentManualCompAxes); + +mlir::WalkResult setManualAxes( + Operation* op, ManualAxesAttr manualAxes, Attribute meshOrRef, + const mlir::SymbolTable& symbolTable, + ManualComputationToParentManualAxes& parentManualCompAxes) { + if (auto manualCompOp = mlir::dyn_cast(op)) { + // Record parent manual axes for this manualCompOp and process its body. + SmallVector& parentAxes = parentManualCompAxes[manualCompOp]; + parentAxes.assign(manualAxes.getValue().begin(), + manualAxes.getValue().end()); + setManualAxesForOpsInBody(manualCompOp, symbolTable, parentManualCompAxes); return mlir::WalkResult::skip(); } if (!mlir::isa(op)) { @@ -202,14 +211,16 @@ mlir::WalkResult setManualAxes(Operation* op, ManualAxesAttr manualAxes, if (CallOp callOp = mlir::dyn_cast(op)) { FuncOp funcOp = symbolTable.lookup(callOp.getCallee()); CHECK(funcOp) << "Failed to lookup function: " << callOp.getCallee().str(); - setFuncManualAxesRecursively(funcOp, manualAxes, meshOrRef, symbolTable); + setFuncManualAxesRecursively(funcOp, manualAxes, meshOrRef, symbolTable, + parentManualCompAxes); } return mlir::WalkResult::advance(); } -void setFuncManualAxesRecursively(FuncOp funcOp, ManualAxesAttr manualAxes, - Attribute meshOrRef, - const mlir::SymbolTable& symbolTable) { +void setFuncManualAxesRecursively( + FuncOp funcOp, ManualAxesAttr manualAxes, Attribute meshOrRef, + const mlir::SymbolTable& symbolTable, + ManualComputationToParentManualAxes& parentManualCompAxes) { llvm::SmallVector funcArgAttrs; funcArgAttrs.reserve(funcOp.getNumArguments()); for (int argNum = 0; argNum < funcOp.getNumArguments(); argNum++) { @@ -250,15 +261,15 @@ void setFuncManualAxesRecursively(FuncOp funcOp, ManualAxesAttr manualAxes, // Walk in preorder of blocks in order to stop walks on manual computations. funcOp->walk([&](Operation* op) { - return setManualAxes(op, manualAxes, meshOrRef, symbolTable); + return setManualAxes(op, manualAxes, meshOrRef, symbolTable, + parentManualCompAxes); }); } // Sets the manual axes of all operations in `op`'s body. void setManualAxesForOpsInBody( - ManualComputationOp op, - const ManualComputationToParentManualAxes& parentManualCompAxes, - const mlir::SymbolTable& symbolTable) { + ManualComputationOp op, const mlir::SymbolTable& symbolTable, + ManualComputationToParentManualAxes& parentManualCompAxes) { TensorShardingAttr sharding = getFirstSharding(op); if (!sharding) { // If there are no in/out shardings, op.getManualAxes() must be empty. We do @@ -280,7 +291,8 @@ void setManualAxesForOpsInBody( // Set the manual axes of all operations in the body. op.getBody().front().walk( [&](Operation* opInBody) { - return setManualAxes(opInBody, manualAxesAttr, meshOrRef, symbolTable); + return setManualAxes(opInBody, manualAxesAttr, meshOrRef, symbolTable, + parentManualCompAxes); }); } @@ -488,7 +500,7 @@ class ShardMapExportPass parentAxes.insert(parentAxes.end(), parentOp.getManualAxes().begin(), parentOp.getManualAxes().end()); } - setManualAxesForOpsInBody(op, parentManualCompAxes, symbolTable); + setManualAxesForOpsInBody(op, symbolTable, parentManualCompAxes); }); // Need to do a separate post order walk to inline the diff --git a/third_party/xla/xla/service/spmd/shardy/test/stablehlo_round_trip_shard_map_export.mlir b/third_party/xla/xla/service/spmd/shardy/test/stablehlo_round_trip_shard_map_export.mlir index 8a7d5aa3cbd6b4..5ca34525922de5 100644 --- a/third_party/xla/xla/service/spmd/shardy/test/stablehlo_round_trip_shard_map_export.mlir +++ b/third_party/xla/xla/service/spmd/shardy/test/stablehlo_round_trip_shard_map_export.mlir @@ -575,3 +575,43 @@ func.func private @bar(%arg0: tensor<8xi32>) -> tensor<8xi32> { // CHECK-NEXT: %1 = call @bar(%0) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}]>]>, xla.sdy.manual_axes = #sdy} : (tensor<8xi32>) -> tensor<8xi32> // CHECK-NEXT: return %arg0 : tensor<4xi32> // CHECK-NEXT: } + +// ----- +sdy.mesh @mesh = <["a"=2, "b"=4]> + +// Tests that when a function containing a `ManualComputationOp` (callee) is +// called from within another `ManualComputationOp` (caller), the callee's +// `ManualComputationOp` correctly inherits parent manual axes from the caller. +// +// 1. Caller passes parent manual axes {"a"} to callee's inputs/outputs: +// CHECK-LABEL: func private @called_func_with_manual_comp( +// CHECK-SAME: %arg0: tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{}, {}]>, xla.sdy.manual_axes = #sdy}) +// CHECK-SAME: -> (tensor<8x8xf32> {sdy.sharding = #sdy.sharding<@mesh, [{}, {}]>, xla.sdy.manual_axes = #sdy}) { +// CHECK-NEXT: %[[COPY:.*]] = mhlo.copy %arg0 {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}, {"b"}]>]>, xla.sdy.manual_axes = #sdy} : tensor<8x8xf32> +// +// 2. Full-to-shard inside callee inherits caller axis "a" + callee axis "b" -> manual axes {"a", "b"}: +// CHECK-NEXT: %[[FULL_TO_SHARD:.*]] = stablehlo.custom_call @SPMDFullToShardShape(%[[COPY]]) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}, {}]>]>, xla.sdy.manual_axes = #sdy} : (tensor<8x8xf32>) -> tensor<8x2xf32> +// +// 3. Inner body call runs with both axes manual: +// CHECK-NEXT: %[[CALL:.*]] = call @xla.sdy.inlinable_manual_computation_body(%[[FULL_TO_SHARD]]) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}, {}]>]>, xla.sdy.manual_axes = #sdy} : (tensor<8x2xf32>) -> tensor<8x2xf32> +// CHECK-NEXT: %[[COPY_1:.*]] = mhlo.copy %[[CALL]] {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}, {}]>]>, xla.sdy.manual_axes = #sdy} : tensor<8x2xf32> +// +// 4. Shard-to-full returns to caller scope, with manual axis {"a"} preserved: +// CHECK-NEXT: %[[SHARD_TO_FULL:.*]] = stablehlo.custom_call @SPMDShardToFullShape(%[[COPY_1]]) {sdy.sharding = #sdy.sharding_per_value<[<@mesh, [{}, {"b"}]>]>, xla.sdy.manual_axes = #sdy} : (tensor<8x2xf32>) -> tensor<8x8xf32> +// CHECK-NEXT: return %[[SHARD_TO_FULL]] : tensor<8x8xf32> +func.func private @called_func_with_manual_comp(%arg0: tensor<8x8xf32>) -> tensor<8x8xf32> { + %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{}, {"b"}]>] out_shardings=[<@mesh, [{}, {"b"}]>] manual_axes={"b"} (%arg1: tensor<8x2xf32>) { + %1 = stablehlo.add %arg1, %arg1 : tensor<8x2xf32> + sdy.return %1 : tensor<8x2xf32> + } : (tensor<8x8xf32>) -> tensor<8x8xf32> + return %0 : tensor<8x8xf32> +} + +// CHECK-LABEL: func @manual_comp_calls_func_with_manual_comp( +func.func @manual_comp_calls_func_with_manual_comp(%arg0: tensor<16x8xf32>) -> tensor<16x8xf32> { + %0 = sdy.manual_computation(%arg0) in_shardings=[<@mesh, [{"a"}, {}]>] out_shardings=[<@mesh, [{"a"}, {}]>] manual_axes={"a"} (%arg1: tensor<8x8xf32>) { + %1 = func.call @called_func_with_manual_comp(%arg1) : (tensor<8x8xf32>) -> tensor<8x8xf32> + sdy.return %1 : tensor<8x8xf32> + } : (tensor<16x8xf32>) -> tensor<16x8xf32> + return %0 : tensor<16x8xf32> +} From 9c6fd2643ea9b23f0391eea31a60cde5800305a5 Mon Sep 17 00:00:00 2001 From: Bhatu Date: Tue, 1 Sep 2026 15:34:53 -0700 Subject: [PATCH 43/43] Propagate min/max reduction identity element constraints in ConstraintPropagator. PiperOrigin-RevId: 974744064 --- .../xla/xla/tests/constraint_propagator.cc | 36 +++++++-- .../xla/xla/tests/constraint_propagator.h | 11 ++- .../xla/tests/constraint_propagator_test.cc | 77 +++++++++++++++++++ third_party/xla/xla/tests/test_utils.cc | 4 + 4 files changed, 121 insertions(+), 7 deletions(-) diff --git a/third_party/xla/xla/tests/constraint_propagator.cc b/third_party/xla/xla/tests/constraint_propagator.cc index b504488f97fbbe..37cf1f89250ad2 100644 --- a/third_party/xla/xla/tests/constraint_propagator.cc +++ b/third_party/xla/xla/tests/constraint_propagator.cc @@ -45,7 +45,6 @@ namespace xla { IdentityElementType GetReductionIdentityElementType( const HloComputation& computation) { - // TODO(b/77635120): Add init values, for min, max, and their arg variants. const HloInstruction* const root = computation.root_instruction(); if (computation.num_parameters() != 2 || root->operand_count() != 2 || root->operand(0)->opcode() != HloOpcode::kParameter || @@ -56,9 +55,15 @@ IdentityElementType GetReductionIdentityElementType( switch (root->opcode()) { case HloOpcode::kAdd: + case HloOpcode::kOr: return IdentityElementType::kZero; case HloOpcode::kMultiply: + case HloOpcode::kAnd: return IdentityElementType::kOne; + case HloOpcode::kMaximum: + return IdentityElementType::kMinimum; + case HloOpcode::kMinimum: + return IdentityElementType::kMaximum; default: return IdentityElementType::kUnknown; } @@ -657,19 +662,29 @@ absl::Status ConstraintPropagator::SeedConstraints( case HloOpcode::kReduce: case HloOpcode::kReduceWindow: { - int64_t first_init = inst->operand_count() / 2; - if (inst->opcode() == HloOpcode::kReduceWindow) { - first_init = 1; - } + int64_t first_init = inst->opcode() == HloOpcode::kReduce + ? inst->operand_count() / 2 + : inst->operand_count() - 1; IdentityElementType etype = GetReductionIdentityElementType(*inst->to_apply()); for (int64_t i = first_init; i < inst->operand_count(); ++i) { + PrimitiveType elem_type = inst->operand(i)->shape().element_type(); if (etype == IdentityElementType::kZero) { states_[inst->operand(i)].AddConstraint( ConstraintInterval{0.0, 0.0, false}); } else if (etype == IdentityElementType::kOne) { states_[inst->operand(i)].AddConstraint( ConstraintInterval{1.0, 1.0, true}); + } else if (etype == IdentityElementType::kMinimum) { + if (auto domain = GetTypeFiniteDomain(elem_type)) { + states_[inst->operand(i)].AddConstraint( + ConstraintInterval{domain->min, domain->min, false}); + } + } else if (etype == IdentityElementType::kMaximum) { + if (auto domain = GetTypeFiniteDomain(elem_type)) { + states_[inst->operand(i)].AddConstraint( + ConstraintInterval{domain->max, domain->max, false}); + } } } break; @@ -678,12 +693,23 @@ absl::Status ConstraintPropagator::SeedConstraints( case HloOpcode::kSelectAndScatter: { IdentityElementType etype = GetReductionIdentityElementType(*inst->scatter()); + PrimitiveType elem_type = inst->operand(2)->shape().element_type(); if (etype == IdentityElementType::kZero) { states_[inst->operand(2)].AddConstraint( ConstraintInterval{0.0, 0.0, false}); } else if (etype == IdentityElementType::kOne) { states_[inst->operand(2)].AddConstraint( ConstraintInterval{1.0, 1.0, true}); + } else if (etype == IdentityElementType::kMinimum) { + if (auto domain = GetTypeFiniteDomain(elem_type)) { + states_[inst->operand(2)].AddConstraint( + ConstraintInterval{domain->min, domain->min, false}); + } + } else if (etype == IdentityElementType::kMaximum) { + if (auto domain = GetTypeFiniteDomain(elem_type)) { + states_[inst->operand(2)].AddConstraint( + ConstraintInterval{domain->max, domain->max, false}); + } } break; } diff --git a/third_party/xla/xla/tests/constraint_propagator.h b/third_party/xla/xla/tests/constraint_propagator.h index f78f34369598c5..9fa12d6fb24c36 100644 --- a/third_party/xla/xla/tests/constraint_propagator.h +++ b/third_party/xla/xla/tests/constraint_propagator.h @@ -30,10 +30,17 @@ limitations under the License. namespace xla { -enum class IdentityElementType { kUnknown, kZero, kOne }; +enum class IdentityElementType { + kUnknown, + kZero, + kOne, + kMinimum, + kMaximum, +}; // Returns the identity element type for the given reduction computation. -// Add => 0, Mul => 1, etc.. Returns kUnknown otherwise. +// Add => 0, Mul => 1, Max => MinValue, Min => MaxValue. Returns kUnknown +// otherwise. IdentityElementType GetReductionIdentityElementType( const HloComputation& computation); diff --git a/third_party/xla/xla/tests/constraint_propagator_test.cc b/third_party/xla/xla/tests/constraint_propagator_test.cc index 04c53c8656f156..651fdc2ecb8a0e 100644 --- a/third_party/xla/xla/tests/constraint_propagator_test.cc +++ b/third_party/xla/xla/tests/constraint_propagator_test.cc @@ -1394,5 +1394,82 @@ ENTRY main { EXPECT_DOUBLE_EQ(indices_int.min, 0.0); EXPECT_DOUBLE_EQ(indices_int.max, 31.0); } + +TEST_F(ConstraintPropagatorTest, ReduceMaxIdentityElementConstraint) { + const char* hlo = R"( +HloModule TestModule +max_computation { + x = bf16[] parameter(0) + y = bf16[] parameter(1) + ROOT max = bf16[] maximum(x, y) +} +ENTRY main { + param_0 = bf16[4,5,128,256] parameter(0) + init_val = bf16[] parameter(1) + ROOT reduce = bf16[4,5] reduce(param_0, init_val), dimensions={2,3}, + to_apply=max_computation +} +)"; + ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(hlo)); + ASSERT_OK_AND_ASSIGN(auto states, ConstraintPropagator::Run(*module)); + + auto init_int = states[module->entry_computation()->parameter_instruction(1)] + .GetConstraintInterval(); + EXPECT_FALSE(init_int.IsEmpty()); + EXPECT_EQ(init_int.min, -65504.0); + EXPECT_EQ(init_int.max, -65504.0); +} + +TEST_F(ConstraintPropagatorTest, ReduceMinIdentityElementConstraint) { + const char* hlo = R"( +HloModule TestModule +min_computation { + x = bf16[] parameter(0) + y = bf16[] parameter(1) + ROOT min = bf16[] minimum(x, y) +} +ENTRY main { + param_0 = bf16[4,5,128,256] parameter(0) + init_val = bf16[] parameter(1) + ROOT reduce = bf16[4,5] reduce(param_0, init_val), dimensions={2,3}, + to_apply=min_computation +} +)"; + ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(hlo)); + ASSERT_OK_AND_ASSIGN(auto states, ConstraintPropagator::Run(*module)); + + auto init_int = states[module->entry_computation()->parameter_instruction(1)] + .GetConstraintInterval(); + EXPECT_FALSE(init_int.IsEmpty()); + EXPECT_EQ(init_int.min, 65504.0); + EXPECT_EQ(init_int.max, 65504.0); +} + +TEST_F(ConstraintPropagatorTest, GetReductionIdentityElementTypeMaxAndMin) { + const char* hlo = R"( +HloModule TestModule +max_computation { + x = f32[] parameter(0) + y = f32[] parameter(1) + ROOT max = f32[] maximum(x, y) +} +min_computation { + x = f32[] parameter(0) + y = f32[] parameter(1) + ROOT min = f32[] minimum(x, y) +} +ENTRY main { + x = f32[] parameter(0) + ROOT root = f32[] negate(x) +} +)"; + ASSERT_OK_AND_ASSIGN(auto module, ParseAndReturnVerifiedModule(hlo)); + EXPECT_EQ(GetReductionIdentityElementType( + *module->GetComputationWithName("max_computation")), + IdentityElementType::kMinimum); + EXPECT_EQ(GetReductionIdentityElementType( + *module->GetComputationWithName("min_computation")), + IdentityElementType::kMaximum); +} } // namespace } // namespace xla diff --git a/third_party/xla/xla/tests/test_utils.cc b/third_party/xla/xla/tests/test_utils.cc index 49135ab69658a7..c846b0df2734d3 100644 --- a/third_party/xla/xla/tests/test_utils.cc +++ b/third_party/xla/xla/tests/test_utils.cc @@ -311,6 +311,10 @@ absl::StatusOr CreateLiteralForConstrainedUses( return LiteralUtil::Zero(param_shape.element_type()); case IdentityElementType::kOne: return LiteralUtil::One(param_shape.element_type()); + case IdentityElementType::kMinimum: + return LiteralUtil::MinValue(param_shape.element_type()); + case IdentityElementType::kMaximum: + return LiteralUtil::MaxValue(param_shape.element_type()); case IdentityElementType::kUnknown: // We want the identity element for the computation, but we don't // really know what it is - so any value we generate will be just as