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. 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: diff --git a/tensorflow/core/kernels/batch_kernels.cc b/tensorflow/core/kernels/batch_kernels.cc index 6223aa1af4f547..58aadb27d12e17 100644 --- a/tensorflow/core/kernels/batch_kernels.cc +++ b/tensorflow/core/kernels/batch_kernels.cc @@ -860,6 +860,20 @@ 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 (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 " @@ -1135,6 +1149,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/core/kernels/maxpooling_op.cc b/tensorflow/core/kernels/maxpooling_op.cc index bbcbbe8269649c..6f7d68cfdec580 100644 --- a/tensorflow/core/kernels/maxpooling_op.cc +++ b/tensorflow/core/kernels/maxpooling_op.cc @@ -1100,9 +1100,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 86ca4f756bea86..eab784a8623739 100644 --- a/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc +++ b/tensorflow/core/kernels/maxpooling_op_gpu.cu.cc @@ -208,11 +208,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]); + } } } @@ -443,7 +447,7 @@ absl::Status 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() ? absl::OkStatus() : absl::InternalError("GPU execution failed"); } diff --git a/tensorflow/core/kernels/quantize_and_dequantize_op.cc b/tensorflow/core/kernels/quantize_and_dequantize_op.cc index 2ce48048465534..b7451da6c95be4 100644 --- a/tensorflow/core/kernels/quantize_and_dequantize_op.cc +++ b/tensorflow/core/kernels/quantize_and_dequantize_op.cc @@ -305,6 +305,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/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..fdf4da61fe6d50 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,22 @@ TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitRankNotOne) { "Ragged splits must have rank 1")); } +TEST_F(RaggedTensorFromVariantKernelTest, RaggedSplitEmpty) { + 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}), + {variant_component_1}); + 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}; 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/core/kernels/rnn/lstm_ops.cc b/tensorflow/core/kernels/rnn/lstm_ops.cc index 3609541988cccd..6307a62c91d3d0 100644 --- a/tensorflow/core/kernels/rnn/lstm_ops.cc +++ b/tensorflow/core/kernels/rnn/lstm_ops.cc @@ -1058,13 +1058,10 @@ 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, + OP_REQUIRES(ctx, seq_len_max >= 0 && seq_len_max <= timelen, absl::InvalidArgumentError( - absl::StrCat("seq_len_max must be >= 0: ", seq_len_max))); - OP_REQUIRES(ctx, seq_len_max <= timelen, - absl::InvalidArgumentError(absl::StrCat( - "seq_len_max must be <= timelen (x.dim_size(0)): ", - seq_len_max, " vs. ", timelen))); + 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"); @@ -1389,13 +1386,10 @@ 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, + OP_REQUIRES(ctx, seq_len_max >= 0 && seq_len_max <= timelen, absl::InvalidArgumentError( - absl::StrCat("seq_len_max must be >= 0: ", seq_len_max))); - OP_REQUIRES(ctx, seq_len_max <= timelen, - absl::InvalidArgumentError(absl::StrCat( - "seq_len_max must be <= timelen (x.dim_size(0)): ", - seq_len_max, " vs. ", timelen))); + 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"); diff --git a/tensorflow/core/kernels/scatter_nd_op.cc b/tensorflow/core/kernels/scatter_nd_op.cc index a1c8bdd66d15b8..7fc4321eb642c0 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/core/kernels/string_ngrams_op.cc b/tensorflow/core/kernels/string_ngrams_op.cc index 75f9d2c25aed66..80e7e052afae1c 100644 --- a/tensorflow/core/kernels/string_ngrams_op.cc +++ b/tensorflow/core/kernels/string_ngrams_op.cc @@ -78,10 +78,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/core/kernels/unicode_ops.cc b/tensorflow/core/kernels/unicode_ops.cc index b75e2f41e56230..79b068422bf1d8 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, @@ -453,9 +453,9 @@ class UnicodeDecodeBaseOp : public OpKernel { Tensor* output_char_values; OP_REQUIRES_OK( - ctx, ctx->allocate_output( - "char_values", {static_cast(char_values.size())}, - &output_char_values)); + ctx, ctx->allocate_output("char_values", + {static_cast(char_values.size())}, + &output_char_values)); auto out_char_values = output_char_values->vec(); if (generate_offsets_) { DCHECK(offset_values.size() == char_values.size()); @@ -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) { 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(); diff --git a/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin b/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin index 98c198ee6baf62..031f8772277612 100644 Binary files a/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin and b/tensorflow/lite/experimental/acceleration/compatibility/gpu_compatibility.bin differ 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 be29ce381f750e..5298c63077fca5 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,45 @@ 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): 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..e71d6d714f1718 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,24 @@ 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 + # 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]) + # 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|Must have updates\.shape", + ): + 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 +848,24 @@ 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) + # 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|Must have updates\.shape", + ): + 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): 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 a6c7e3e2fa0297..c7f707285e1e23 100644 --- a/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/pooling_ops_test.py @@ -2618,6 +2618,29 @@ 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(): 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..9309f178ee9122 100644 --- a/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py +++ b/tensorflow/python/kernel_tests/nn_ops/rnn_cell_test.py @@ -1468,6 +1468,181 @@ 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): 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 043cee80921463..40d8487f560a01 100644 --- a/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py +++ b/tensorflow/python/kernel_tests/quantization_ops/quantization_ops_test.py @@ -510,6 +510,34 @@ 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|Shape must be rank 0)", + ): + 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() 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..ee14f88eb1cae3 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,24 @@ 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]) + 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 class UnicodeSplitTest(test_util.TensorFlowTestCase, diff --git a/tensorflow/python/ops/batch_ops_test.py b/tensorflow/python/ops/batch_ops_test.py index 15a1a71a93bb63..698d25ded44347 100644 --- a/tensorflow/python/ops/batch_ops_test.py +++ b/tensorflow/python/ops/batch_ops_test.py @@ -242,7 +242,8 @@ 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) @@ -254,7 +255,64 @@ def testUnbatchInvalidIdArg(self): id=arg_id, timeout_micros=50, container="", - shared_name="") + 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 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.""" diff --git a/tensorflow/python/ops/raw_ops_test.py b/tensorflow/python/ops/raw_ops_test.py index a6f19007b5d7e8..e1de9ac5fe6632 100644 --- a/tensorflow/python/ops/raw_ops_test.py +++ b/tensorflow/python/ops/raw_ops_test.py @@ -76,6 +76,40 @@ 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 + ): + 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"] data_splits = [0, 1] 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)) 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/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_ 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/debug_options_flags.cc b/third_party/xla/xla/debug_options_flags.cc index 540880d59af90f..8d36c0bf83f141 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( @@ -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/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> } 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, 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 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)); 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."; 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()), 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()) { 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( 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( 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> +} diff --git a/third_party/xla/xla/tests/constraint_propagator.cc b/third_party/xla/xla/tests/constraint_propagator.cc index 70b33e7bd1000a..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; } @@ -394,13 +399,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 @@ -663,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; @@ -684,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 da5540e06f576c..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 @@ -422,31 +426,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 +466,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 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 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.