<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/82382>82382</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
MLIR parallel-loops-to-gpu pass crash
</td>
</tr>
<tr>
<th>Labels</th>
<td>
mlir
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
imesmuaran
</td>
</tr>
</table>
<pre>
While lowering some MLIR code to test which parameters I can tune for more performance, I encountered a crash.
The crash can be reproduced with the following minimal example:
```
func.func @main() {
%in_buf = memref.alloc() : memref<16x230x230x3xf32>
%filter_buf = memref.alloc() : memref<64x7x7x3xf32>
%out_buf = memref.alloc() : memref<16x112x112x64xf32>
linalg.conv_2d_nhwc_fhwc {dilations = dense<1> : tensor<2xi64>, strides = dense<2> : tensor<2xi64>}
ins (%in_buf, %filter_buf: memref<16x230x230x3xf32>, memref<64x7x7x3xf32>)
outs (%out_buf: memref<16x112x112x64xf32>)
return
}
```
while trying to generate a `gpu.launch` op using the following pass pipeline:
```
mlir-opt --convert-linalg-to-affine-loops \
--affine-loop-tile="tile-sizes=4,28,28,16,7,7,3" \
--affine-loop-unroll="unroll-factor=4" \
--canonicalize \
--affine-parallelize \
--lower-affine \
--canonicalize \
--gpu-map-parallel-loops \
--convert-parallel-loops-to-gpu \
conv2d.mlir
```
The crash is not produced when the `tile-sizes` is tweaked e.g. replacing one of the `7` with a `6`.
# Backtrace
```
PLEASE submit a bug report to https://github.com/llvm/llvm-project/issues/ and include the crash backtrace.
Stack dump:
0. Program arguments: mlir-opt --convert-linalg-to-affine-loops --affine-loop-tile=tile-sizes=4,28,28,16,7,7,3 --affine-loop-unroll=unroll-factor=4 --canonicalize --affine-parallelize --lower-affine --canonicalize --gpu-map-parallel-loops --convert-parallel-loops-to-gpu conv2d.mlir
#0 0x00005c544a6cdee6 llvm::sys::PrintStackTrace(llvm::raw_ostream&, int) /home/j/llvm-project/llvm/lib/Support/Unix/Signals.inc:723:13
#1 0x00005c544a6cbe40 llvm::sys::RunSignalHandlers() /home/j/llvm-project/llvm/lib/Support/Signals.cpp:106:18
#2 0x00005c544a6ce88b SignalHandler(int) /home/j/llvm-project/llvm/lib/Support/Unix/Signals.inc:413:1
#3 0x000073b19aa5a770 (/usr/lib/libc.so.6+0x3c770)
#4 0x00005c544d35cf45 void mlir::detail::IROperandBase::insertInto<mlir::IRObjectWithUseList<mlir::OpOperand>>(mlir::IRObjectWithUseList<mlir::OpOperand>*) /home/j/llvm-project/mlir/include/mlir/IR/UseDefLists.h:99:24
#5 0x00005c544d35cf45 mlir::IROperand<mlir::OpOperand, mlir::Value>::insertIntoCurrent() /home/j/llvm-project/mlir/include/mlir/IR/UseDefLists.h:186:30
#6 0x00005c544d35cf45 mlir::IROperand<mlir::OpOperand, mlir::Value>::IROperand(mlir::Operation*, mlir::Value) /home/j/llvm-project/mlir/include/mlir/IR/UseDefLists.h:132:5
#7 0x00005c544d35cf45 mlir::OpOperand::OpOperand(mlir::Operation*, mlir::Value) /home/j/llvm-project/mlir/include/mlir/IR/Value.h:280:38
#8 0x00005c544d35cf45 mlir::detail::OperandStorage::OperandStorage(mlir::Operation*, mlir::OpOperand*, mlir::ValueRange) /home/j/llvm-project/mlir/lib/IR/OperationSupport.cpp:245:30
#9 0x00005c544d351d33 mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::DictionaryAttr, mlir::OpaqueProperties, mlir::BlockRange, unsigned int) /home/j/llvm-project/mlir/lib/IR/Operation.cpp:122:3
#10 0x00005c544d351565 mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::NamedAttrList&&, mlir::OpaqueProperties, mlir::BlockRange, unsigned int) /home/j/llvm-project/mlir/lib/IR/Operation.cpp:75:10
#11 0x00005c544d351565 mlir::Operation::create(mlir::Location, mlir::OperationName, mlir::TypeRange, mlir::ValueRange, mlir::NamedAttrList&&, mlir::OpaqueProperties, mlir::BlockRange, mlir::RegionRange) /home/j/llvm-project/mlir/lib/IR/Operation.cpp:58:7
#12 0x00005c544d351407 mlir::Operation::create(mlir::OperationState const&) /home/j/llvm-project/mlir/lib/IR/Operation.cpp:36:7
#13 0x00005c544d2a6db0 mlir::OpBuilder::create(mlir::OperationState const&) /home/j/llvm-project/mlir/lib/IR/Builders.cpp:447:17
#14 0x00005c544cb5558b llvm::ValueIsPresent<mlir::Operation*, void>::isPresent(mlir::Operation* const&) /home/j/llvm-project/llvm/include/llvm/Support/Casting.h:622:55
#15 0x00005c544cb5558b bool llvm::detail::isPresent<mlir::Operation*>(mlir::Operation* const&) /home/j/llvm-project/llvm/include/llvm/Support/Casting.h:630:10
#16 0x00005c544cb5558b decltype(auto) llvm::dyn_cast<mlir::affine::AffineApplyOp, mlir::Operation>(mlir::Operation*) /home/j/llvm-project/llvm/include/llvm/Support/Casting.h:662:3
#17 0x00005c544cb5558b mlir::affine::AffineApplyOp mlir::OpBuilder::create<mlir::affine::AffineApplyOp, mlir::AffineMap, mlir::ValueRange>(mlir::Location, mlir::AffineMap&&, mlir::ValueRange&&) /home/j/llvm-project/mlir/include/mlir/IR/Builders.h:496:19
#18 0x00005c544cb547a3 mlir::Operation::getOpResultImpl(unsigned int) /home/j/llvm-project/mlir/include/mlir/IR/Operation.h:1000:5
#19 0x00005c544cb547a3 mlir::Operation::getResult(unsigned int) /home/j/llvm-project/mlir/include/mlir/IR/Operation.h:402:54
#20 0x00005c544cb547a3 mlir::OpTrait::OneTypedResult<mlir::IndexType>::Impl<mlir::affine::AffineApplyOp>::getResult() /home/j/llvm-project/mlir/include/mlir/IR/OpDefinition.h:697:33
#21 0x00005c544cb547a3 mlir::OpTrait::OneTypedResult<mlir::IndexType>::Impl<mlir::affine::AffineApplyOp>::operator mlir::detail::TypedValue<mlir::IndexType>() /home/j/llvm-project/mlir/include/mlir/IR/OpDefinition.h:702:54
#22 0x00005c544cb547a3 processParallelLoop(mlir::scf::ParallelOp, mlir::gpu::LaunchOp, mlir::IRMapping&, llvm::SmallVectorImpl<mlir::Operation*>&, llvm::DenseMap<mlir::gpu::Processor, mlir::Value, llvm::DenseMapInfo<mlir::gpu::Processor, void>, llvm::detail::DenseMapPair<mlir::gpu::Processor, mlir::Value>>&, mlir::PatternRewriter&) /home/j/llvm-project/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp:457:18
#23 0x00005c544cb53b93 mlir::LogicalResult::failed() const /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#24 0x00005c544cb53b93 mlir::failed(mlir::LogicalResult) /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#25 0x00005c544cb53b93 (anonymous namespace)::ParallelToGpuLaunchLowering::matchAndRewrite(mlir::scf::ParallelOp, mlir::PatternRewriter&) const /home/j/llvm-project/mlir/lib/Conversion/SCFToGPU/SCFToGPU.cpp:642:11
#26 0x00005c544f524be1 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<mlir::LogicalResult (mlir::Pattern const&)>)::$_0::operator()() const /home/j/llvm-project/mlir/lib/Rewrite/PatternApplicator.cpp:208:13
#27 0x00005c544f524be1 void llvm::function_ref<void ()>::callback_fn<mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<mlir::LogicalResult (mlir::Pattern const&)>)::$_0>(long) /home/j/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#28 0x00005c544f521910 mlir::PatternApplicator::matchAndRewrite(mlir::Operation*, mlir::PatternRewriter&, llvm::function_ref<bool (mlir::Pattern const&)>, llvm::function_ref<void (mlir::Pattern const&)>, llvm::function_ref<mlir::LogicalResult (mlir::Pattern const&)>) /home/j/llvm-project/mlir/lib/Rewrite/PatternApplicator.cpp:229:9
#29 0x00005c544d23afdd (anonymous namespace)::OperationLegalizer::legalize(mlir::Operation*, mlir::ConversionPatternRewriter&) /home/j/llvm-project/mlir/lib/Transforms/Utils/DialectConversion.cpp:0:0
#30 0x00005c544d2318f9 (anonymous namespace)::OperationConverter::convert(mlir::ConversionPatternRewriter&, mlir::Operation*) /home/j/llvm-project/mlir/lib/Transforms/Utils/DialectConversion.cpp:0:0
#31 0x00005c544d2318f9 (anonymous namespace)::OperationConverter::convertOperations(llvm::ArrayRef<mlir::Operation*>, llvm::function_ref<void (mlir::Diagnostic&)>) /home/j/llvm-project/mlir/lib/Transforms/Utils/DialectConversion.cpp:2433:16
#32 0x00005c544d2346d4 mlir::applyPartialConversion(llvm::ArrayRef<mlir::Operation*>, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, llvm::DenseSet<mlir::Operation*, llvm::DenseMapInfo<mlir::Operation*, void>>*) /home/j/llvm-project/mlir/lib/Transforms/Utils/DialectConversion.cpp:3403:22
#33 0x00005c544d2346d4 mlir::applyPartialConversion(mlir::Operation*, mlir::ConversionTarget const&, mlir::FrozenRewritePatternSet const&, llvm::DenseSet<mlir::Operation*, llvm::DenseMapInfo<mlir::Operation*, void>>*) /home/j/llvm-project/mlir/lib/Transforms/Utils/DialectConversion.cpp:3409:10
#34 0x00005c544cb51f0a (anonymous namespace)::ParallelLoopToGpuPass::runOnOperation() /home/j/llvm-project/mlir/lib/Conversion/SCFToGPU/SCFToGPUPass.cpp:62:16
#35 0x00005c544d1d301f mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_1::operator()() const /home/j/llvm-project/mlir/lib/Pass/Pass.cpp:0:17
#36 0x00005c544d1d301f void llvm::function_ref<void ()>::callback_fn<mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int)::$_1>(long) /home/j/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#37 0x00005c544d1d301f llvm::function_ref<void ()>::operator()() const /home/j/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#38 0x00005c544d1d301f void mlir::MLIRContext::executeAction<mlir::PassExecutionAction, mlir::Pass&>(llvm::function_ref<void ()>, llvm::ArrayRef<mlir::IRUnit>, mlir::Pass&) /home/j/llvm-project/mlir/include/mlir/IR/MLIRContext.h:275:7
#39 0x00005c544d1d301f mlir::detail::OpToOpPassAdaptor::run(mlir::Pass*, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int) /home/j/llvm-project/mlir/lib/Pass/Pass.cpp:513:21
#40 0x00005c544d1d5c15 mlir::LogicalResult::failed() const /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#41 0x00005c544d1d5c15 mlir::failed(mlir::LogicalResult) /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#42 0x00005c544d1d5c15 mlir::detail::OpToOpPassAdaptor::runPipeline(mlir::OpPassManager&, mlir::Operation*, mlir::AnalysisManager, bool, unsigned int, mlir::PassInstrumentor*, mlir::PassInstrumentation::PipelineParentInfo const*) /home/j/llvm-project/mlir/lib/Pass/Pass.cpp:585:9
#43 0x00005c544d1d5c15 mlir::PassManager::runPasses(mlir::Operation*, mlir::AnalysisManager) /home/j/llvm-project/mlir/lib/Pass/Pass.cpp:896:10
#44 0x00005c544d1d5c15 mlir::PassManager::run(mlir::Operation*) /home/j/llvm-project/mlir/lib/Pass/Pass.cpp:876:60
#45 0x00005c544d1c58e9 mlir::LogicalResult::failed() const /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:44:33
#46 0x00005c544d1c58e9 mlir::failed(mlir::LogicalResult) /home/j/llvm-project/mlir/include/mlir/Support/LogicalResult.h:72:58
#47 0x00005c544d1c58e9 performActions(llvm::raw_ostream&, std::shared_ptr<llvm::SourceMgr> const&, mlir::MLIRContext*, mlir::MlirOptMainConfig const&) /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:394:7
#48 0x00005c544d1c4a56 processBuffer(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::MlirOptMainConfig const&, mlir::DialectRegistry&, llvm::ThreadPool*) /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:459:12
#49 0x00005c544d1c4a56 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:530:12
#50 0x00005c544d1c4a56 mlir::LogicalResult llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::callback_fn<mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&)::$_0>(long, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) /home/j/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:45:12
#51 0x00005c544d26efab llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>::operator()(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&) const /home/j/llvm-project/llvm/include/llvm/ADT/STLFunctionalExtras.h:68:12
#52 0x00005c544d26efab mlir::splitAndProcessBuffer(std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::function_ref<mlir::LogicalResult (std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, llvm::raw_ostream&)>, llvm::raw_ostream&, bool, bool) /home/j/llvm-project/mlir/lib/Support/ToolUtilities.cpp:28:12
#53 0x00005c544d1c0d47 mlir::MlirOptMain(llvm::raw_ostream&, std::unique_ptr<llvm::MemoryBuffer, std::default_delete<llvm::MemoryBuffer>>, mlir::DialectRegistry&, mlir::MlirOptMainConfig const&) /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:533:10
#54 0x00005c544d1c0f12 mlir::MlirOptMain(int, char**, llvm::StringRef, llvm::StringRef, mlir::DialectRegistry&) /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:568:14
#55 0x00005c544d1c11f5 mlir::MlirOptMain(int, char**, llvm::StringRef, mlir::DialectRegistry&) /home/j/llvm-project/mlir/lib/Tools/mlir-opt/MlirOptMain.cpp:584:10
#56 0x00005c544a624265 main /home/j/llvm-project/mlir/tools/mlir-opt/mlir-opt.cpp:297:33
#57 0x000073b19aa43cd0 (/usr/lib/libc.so.6+0x25cd0)
#58 0x000073b19aa43d8a __libc_start_main (/usr/lib/libc.so.6+0x25d8a)
#59 0x00005c544a623de5 _start (/home/j/llvm-project/build/bin/mlir-opt+0x405dde5)
[1] 114374 segmentation fault (core dumped) mlir-opt --convert-linalg-to-affine-loops --canonicalize --lower-affine
```
# Environment
```
$ git status
On branch release/18.x
$ mlir-opt --version
LLVM (http://llvm.org/):
LLVM version 18.1.0rc
Optimized build with assertions.
$ uname -a
Linux 6.7.4-arch1-1 #1 SMP PREEMPT_DYNAMIC Mon, 05 Feb 2024 22:07:49 +0000 x86_64 GNU/Linux
```
# Build
```
cmake -G Ninja ../llvm -DLLVM_ENABLE_PROJECTS="clang;llvm;lld;lldb;mlir;openmp" \
-DCMAKE_C_COMPILER=clang \
-DCMAKE_CXX_COMPILER=clang++ \
-DLLVM_ENABLE_LLD=ON \
-DLLVM_BUILD_EXAMPLES=ON \
-DLLVM_TARGETS_TO_BUILD="X86;NVPTX;AMDGPU" \
-DCMAKE_BUILD_TYPE=RelWithDebInfo \
-DLLVM_ENABLE_ASSERTIONS=ON \
-DMLIR_ENABLE_CUDA_RUNNER=ON \
-DMLIR_ENABLE_ROCM_RUNNER=ON \
-DMLIR_ENABLE_SPIRV_CPU_RUNNER=ON \
-DMLIR_INCLUDE_INTEGRATION_TESTS=ON \
-DLLVM_INSTALL_UTILS=ON
cmake --build . --target all
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzsXF1z4jiX_jXOjQrKlj-Ai1zw2csuJBQhPf1eUcKWQdNG8ivJk2R-_ZZkG39gCJC8vZmdqeomYEtH5zx6dHSOZAsJQbYU43vDHRju6A4lcsf4PdljsU8QR_Ruw4K3-992JMIgYi-YE7oFgu0xmM-mS-CzAAPJgMRCgpcd8XcgRhztscRcgCnwEQUyoRiEjIM94xjEmIeM7xH1sQGHYAow9VlCJeY4AAj4HIld2zBHhtlPP1c7nF7VwjYYcBxzFiQ-DsALkTsgd0p8FLEXpdueULJHEcCvaB9H2LAzMYZnZv_0zzChflt9AMMx94hQA3YN2ANGZ1Bu3DB7BnQJXW-SEBj2COzxnuOwjaKI-XkVu59dNuyh5b1C29T_7dfQhoY9LksKSSQxv1ia57x2XjuNklgir1HKsqD-7zlHsvIvEaEo2rZ9Rv9Yw2BNdy_-Oty9-AqTgERIEkaFbi_AVGAl17DHuiWJqWDcsIfwlXiOkg6HQEhOAlytAU_X6IxKKvWIakqZkoGvBFbwexd1ODyNI-yV22KJzBvLUH0fvYqE-heOZcJp9vNgV5WA6eeLHliSvynqSga2mGKOJAYIGJ65jZN2hBLq7wzPBCwGidDlKoSPkRAgJjGOCD1J931EeIvFErRaqn8xl620u1uStVAYEopbEWOxAIY7TKuAVvlGSxI1mEYGhOpbS5A_sTDskWPAoUIu_bA8Aw472X_bgPCkuIRyFkWpwPR7K0S-VJRQMssVgVIaUUaJjyLyJ26SqZxOFOHabVVTO62s2KVCt3HS2qP4IPUIGlDCsVpI4bmNk1JZVQ4GbdUDZ4hQeDkiAGUSFD5uh6nucsMzS9B7piopXzD6iQOA29u2cowR8hUnGMWAhXmtjiqsXaVmlWd4ZsXDGtAGA-T_lBz5uFHHxWzcfxoDkWz2RAIENslWtca4VKTdSRkLxTw4MeBkS-Qu2bR9tjfgJIr-yP-0Ys5-x7404IQIkWBhwAlANACE-lGippEDBJtcmUzNJ4n8nyBI9vGB38qC3oKzLUd7gPg22WMqhR63F1O9keAXs_sUoY_YXGdaI2lrRD2qcoKS79HwiHzAgLYJzFfTNE3Xdx0HeX6AsQd0T9l9w-6LN5F-WXBCpcZ-pZkBu0Uhjl7WTEiO0d6AChdAqNQTD5zs2B4bcPL7ccfndCAbA06eklgxyICTZ0pe1QWypSgSbUJ9w-53oG3Yfcsu9LZqem-wYzbpvUxoKuq_EA0izEU-J96gWq6THyvuWaanPruFTrCmE-52N6DSugG7nw2NY2loCi3sTIuOvbF6CLmo0zHTGW2SCH4QG5GN3xas7RlwYL7afqdjHiYyJcYpGxPYrh86LviDkUAPqhTcAEtEovT7dPkYY45oMEACp5cIFZjLKZXMsIdFrenycaMM_Y3I3bPAMyJk5f5jnElSc6uaXrs31oX9d5HWNeEkczvFhelS4S3wCIeqDdHeGXa_1zPsPnQKjNwmjCrK5so0qqiiksPl7yhKsLa4Ct0w4RwrzrxP26uMsbqKvrZZWOP9x6wpqpb7Ul9TgaTuqOPan2uuCtX6bmFt57y1JSLVzPxVFmgRWnfYNVVXlTxN97zy5XGZ6f0kGUdb3HjtMpNKEDQZu0R0e7HFqQvSZh4azPxc5l2h49bo2avZbAW2DZoU1z99jpGsmjZjfmbZsKneA9rj6q3VW5ybdcrc8vUR8ZUcxN_6UvJ6K-jfCV5wFmMuiYp3yncHEfN_HiQmVOfhwUWz6BlA83kKKubbh_DOMutAup77pYBUAgKFoXbvKqLwvgyaHVdP_QWa1t8KzeL6Em8Jox8e9RmsbldhW6AK66g6ZucqVAu3IlUK7TOaGf9hVW2vqqpdURUiL9iYFVUHCYkCzH-BqllLeYTqOB1F1pKulaDO37iu292UomZNoalYcCxUyFGd56vzgooEi2jlUOXUVHK5VVnsW8yJ2YUiDB4iIQnd6qnR087NdQsT3SYTN4xFJTvL8yN539xaDPpL7LLNmp_xmuwKsB_Jt1jxCSWSKSVKVr7RtY9qAXKaU6bf-_p7P46jt8f4hPs5Z_xn2-zVJqpOk8kXmPLu6LsBkPTWHMUnnW4NqGa3XhJz7IjLHty7wgM0x48HZ6CwdXo6We0V4HZr4DoddCac2mL5GC-xSCI53ceRAbu3zKvNihYOVgfqpmkWkbrStHetpqme_0EdHVN7HeegIzTf03HFEZHZD4rVFB5kalZSYxrgV3WzSJ0U2pfxNa9SRuCjVo9wSCg5GO711IxiF2MUWl_IcKY7ifETmZBuO0tMTzX96Yh1jqgCmxCLOfOxEIts4W7GWFzxJsIPs3W4rETdP23jJPM7eo-gfnu6nKM4JnSbOZ1ilnjaoyj6jn3J-BHkR7Ngre4IU6GdWbnSQZNFahPjjalxk5wpDdkFsvLYoyKj3NO5vAVSgq7WLVt2qnnnBZISc7rEL5xIzK-O0IZ6cVZoNCdPw8mKfVs8l77mQZvbKS0rKsLYNcLYm155iM3YlvgoygeVuhQiEuEgI7OOUG6jdDFPVxpJfaBT8wT12LKm5kGnU5rfOu7OKdmBaW5xULIeHWolVfREGX3bs0QAivZYxHqJu1cdciv2LU7S4TXLdr_TAnsk_V2fBhk1rhu4zbS6vNOuIpfnKEAsqwCkElaGLnQ22DrWTnla4iO9ifGexadXkJpMLQ_hMKF6BWWd7rXqoL0iOhNQDrqP_EBNiF6x_qiQE5S9TG7GIgM6a7M6T6UD9Ophmvb4AfrJUR_lS2hmt7Rpojq709TZGqF38cusSaNoFEUb5P9ch7QCzj9k-WSyqHgkYmrSvjHR6o9WygmsZpNMRxSNXyVHWVKg17JgwY9ujR9WzzL_cQaf2r-fO8I1X4qkDlbXyKGNwiB4b3o7dNEMb_UOc2ZClP28sD-LGeijgdKKIypCxvfCgJNnSSL1d0RQhH1ZtJJBYOr_OQDV7exApblh72IAUuHysFiQ_qzYf97K5hWUy7chP26_9fn2H26LyqZ_n3P0tqyNiKOE4arBOCJoS5mQxL9xxFyFHHRsvXPuFejBGnqOFzjlBSeVby4QlwRF5YDrBlSaGLVCfItlyWWUi004-xPnbMu491QrXUuonvD5JdwL8q-Tq77XbK7f0De2Y9ravxV9Y9_YN9f5rn_6oNwHveoatF1P8azQRJdmTzPGYp1BLZDIns3hCX2kJfMuXnu5LOdRDeV5D6yN9OqTG1Zgm1Z4cgN9xR5jJawfoPgQ9fCE1mZ8IRr2zE9Rrk9R9CaImCOKtlgvQ6gQp2HHsogIrU9MH1J1JyWMzOpeUe1xkAyiz0sWvizEvzjorj2IkuF8FcS3EuI2A7xuzYDuSaIUfTGfTZdDRiV-zZao8Cv2E4n7frpyP6z28ljfJYxm92uZg6KBl_XUhUBVfG3jLD1dPlMij-bnvLUPLQuXzE-f6NGPEZQGW--v448-5GRc_bQiLBagHLNmuetb7tdb3XSs82p-idVNB55X8jIWLfLXB6qxkyp6IMfZXOcmWtWH3JQKyfWD3MqvHa0dlAuUNv9y3ReIYypVKJUHZ1dGSg3M7bqVXNuxz4NdxusALRICiwuD0iPoPqZ_N90BLqI5x7negNsfBXhPu47SzitpV4vQfLeLe1_QL3jn1fwafqHTpGT28l06w4p3HugXMnsQVuwQx8E6ltywh6VdRJZwH8-33LDHJ7KncghQ5_o8IvwxlnNE6JDRkGxvfApqxZjOafbZex9q6i1E54lNz6lMvk4tgPEd5Hr5juwgCUP93P5F6CSU_DvBR-jM8Z7xt1xUqXyAQ5REch3gCOvnUk7UydK9SyGrPpWq07sl3hIh-dtRmrracYyChXbK16aTl6DtuL1qvOj0muBuNOwLwt4M54VMrizuN0buv86k08Ben0xeRAQ3faStIIJrnidCdXn9lvX4L4Hmu5nw34XwRWL9Ky05x_NfkN67tRV5D4do8_-AzH9dz_U5qx4ubOrX0nMfcURknwaLWhjxfwHTX5BjZ0sUOWT695qYpYib1aT1LElEJMF5GgLrvVxL8nwzcDp_j1jl86d_u5p_uk4d29CCJ7HNFgn8HUqXA2q7J0-SE7pd4vDM5fPAfL7Bqcsonvx06ymtZYXupxj8qy3rOrWu9KpvPkMHei7YI0IvaloeN5p_zcdl_bljt1N9y9mx_eD9t5yh6wfFW85KTLcuJugisF6ramshEZfrzIr3BAddVBHcqyFiB9gFqchM2mlYNgmJAvVXEaGEysB8dUw3CLBbNOUOLMMdAQCAZTl2xwECbw-LYkC7DtWezzjWhybgQBHiqoMRKqcP1A4nOHOKhQFtMKZ_EM6oUqi5JHTAlkggJJKJSK89UrDhiPo7wHGEkVAgWd32a1W2UzYh3wbU92az73Nl8k5KfUSEPohCIdxmfKt_6bA0O7xDl87qA6vbttom9_Obj7Eke_InDoDukuzcDCEw12sm7bpKCUV7DFooU4TQ5BV47U7baSHu76yWlZ5c8DRfgMVyPJ4vVuvRvx768-kQzNNdDtMFE7wB0IQO0K82mR399gZQvW-aJnjtemvPAd8eng040S1UlDjVEfplkMYy_h79xKD1DTwQ-jsC7XaGFmiNFDbr8UN_MBuvF8vH_x4PV0_pKS1-hOjWsAepPxpEUZB-bgx7kDqjAYsx3ce1Y19Gw3n_f8br4Xr4OF9MZ-OlYY-0rMp5KkW5Hz-OShpwYMBBRWhZz9lsZNijx4e6QF1m8DydjdbjH_35YjZ-qpfLCq36y2_j1dN69ZiWTy3-0fUMe_DwfbH6YdiD_nyk95qbjEsbWf1rMTbs0RJHvxG5G-GNXoxuUipTvP_0NF6upo8PR3rpovPZdJkXHT6P-uvl88ODRuV82eXjcH5p2afFdPl9PVw8n6iQlZ4-DGfPo_F6-rAaf1v2lc7r1fhpdQrQ6cPTqj-brZ9X01lapszOjH-tdIS1Qasl02cwUBQ1EvYuuLeDnt1Dd_je6phdx-larnO3u_e8ILB8PwgDFCJsurbndhzLRtCxrE7Qce_IvRpXJoSm5bmua7c3PQ8FlgcdK7Q8y_UNx8R7RKJ27i7u9Ok0911od-FdhDY4Evp0NAjTqQsa7uiO32v_vUm2wnDMSB85cBAgiYzwvT4frflUFn1kkz7q5i7h0f3t5-doJf83AAD__-gS0WY">