mirror of
https://github.com/openmlsys/openmlsys-zh.git
synced 2026-04-05 19:58:22 +08:00
Update ch06
This commit is contained in:
Corleone
94
chapter_hardware_accelerator/accelerator_architecture.md
Normal file
94
chapter_hardware_accelerator/accelerator_architecture.md
Normal file
@@ -0,0 +1,94 @@
|
||||
加速器基本组成原理
|
||||
------------------
|
||||
|
||||
上节主要介绍了加速器的意义以及设计思路，了解到加速器与通用处理器在设计上的区别，因此加速器的硬件结构与CPU的硬件结构有着根本的不同，通常都是由多种片上缓存以及多种运算单元组成。本章节主要通过GPU的Volta架构作为样例进行介绍。
|
||||
|
||||
### 硬件加速器的架构
|
||||
|
||||
现代GPU在十分有限的面积上实现了极强的计算能力和极高的储存器以及IO带宽。一块高端的GPU中，晶体管数量已经达到主流CPU的两倍，而且显存已经达到了16GB以上,工作频率也达到了1GHz。GPU的体系架构由两部分组成，分别是流处理阵列和存储器系统，两部分通过一个片上互联网络连接。流处理器阵列和存储器系统都可以单独扩展，规格可以根据产品的市场定位单独裁剪。如GV100的组成如 :numref:`gv100`所示：
|
||||
|
||||
![Volta GV100 [@2017NVIDIA] []](../img/ch06/V100.svg)
|
||||
:width:`800px`
|
||||
:label:`gv100`
|
||||
|
||||
- 6个GPU处理集群（GPU Processing Cluster，GPC）, 每个GPC含有：
|
||||
|
||||
- 7个纹理处理集群（Texture Processing Cluster, TPC） (每个TPC含有两个流多处理器（Streaming Multiprocessor, SM）)
|
||||
|
||||
- 14个SM
|
||||
|
||||
- 84个SM, 每个流多处理器含有：
|
||||
|
||||
- 64个32位浮点运算单元
|
||||
|
||||
- 64个32位整数运算单元
|
||||
|
||||
- 32个64位浮点运算单元
|
||||
|
||||
- 8个张量核
|
||||
|
||||
- 4个纹理单元
|
||||
|
||||
- 8个512-bit内存控制器
|
||||
|
||||
一个完整的GV100 GPU含有84个SM，5376个32位浮点运算单元，5376个32位整型运算单元，2688个64位浮点运算单元，672个张量运算单元和336个纹理单元。一对内存控制器控制一个HBM2 DRAM堆栈。图 :numref:`gv100`中展示的为带有84个SM的GV100 GPU(不同的厂商可以使用不同的配置)，Tesla V100则含有80个SM。
|
||||
|
||||
### 硬件加速器的存储单元
|
||||
|
||||
与传统的CPU模型相似，从一个计算机系统主内存DRAM中获取数据的速度相对于处理器的运算速度较慢。对于加速器而言，如果没有缓存进行快速存取，DRAM的带宽非常不足。如果无法快速的在DRAM上获取程序和数据，加速器将因空置而降低利用率。为了缓解DRAM的带宽问题，GPU提供了不同层次的若干区域供程序员存放数据，每块区域的内存都有自己的最大带宽以及延迟。
|
||||
|
||||
- **寄存器文件（Register File）**：片上最快的存储器，但与CPU不同，GPU的每个SM（流多处理器）有上万个寄存器。但当每个线程使用过多的寄存器时，SM中能够调度的线程块数量就会受到限制，可执行的线程总数量会因此受到限制，可执行的线程数量过少会造成硬件无法充分的利用，性能急剧下降。所以要根据算法的需求合理使用寄存器。
|
||||
|
||||
- **共享内存（Shared Memory）**：共享内存实际上是用户可控的一级缓存，每个SM（流多处理器）中有128KB的一级缓存, 开发者可根据应用程序需要配置最大96KB的一级缓存作为共享内存。共享内存的延迟极低，只有几十个时钟周期，几乎与寄存器相当。共享内存具有高达1.5TB/s的带宽，远远高于全局内存的峰值带宽900GB/s。所以说，共享内存的使用对于一个高性能计算工程师来说是一个必须要掌握的一个概念。
|
||||
|
||||
- **全局内存（Global Memory）**：全局内存之所以称为全局，是因为GPU与CPU都可以对它进行读写操作。全局内存对于GPU中的每个线程都是可见的，都可以直接对全局内存进行读写操作。CPU等其他设备可以通过PCI-E总线对其进行读写操作。全局内存也是GPU中容量最大的一块内存，可达16GB之多。同时也是延迟最大的内存，
|
||||
|
||||
- **常量内存（Constant Memory）**：常量内存其实只是全局内存的一种虚拟地址形式，并没有真正的物理硬件内存块。常量内存有两个特性，一个高速缓存，另一个更重要的特性是它支持将某个单个值广播到线程束中的每个线程中。
|
||||
|
||||
- **纹理内存（Texture Memory）**：纹理内存是全局内存的一个特殊形态。当全局内存被绑定为纹理内存时，执行读写操作将通过专用的纹理缓存来加速。在早期的GPU上没有缓存，因此每个SM（流多处理器）上的纹理内存为设备提供了唯一真正缓存数据的方法。然而随着硬件的升级，一级缓存和二级缓存的出现，纹理缓存的这项优势已经荡然无存。纹理内存的另外一个特性，也是最有用的特性就是当访问存储单元时，允许GPU实现硬件相关的操作。比如说使用纹理内存，可以通过归一化的地址对数组进行访问，获取的数据可以通过硬件进行自动插值，从而达到快速处理数据的目的。此外对于二位数组和三维数组，支持硬件级的双线性插值与三线性插值。纹理内存另一个实用的特性是可以根据数组的索引自动处理边界条件，不需要对特殊边缘进行处理即可完成数组内元素操作，从而防止线程中分支的产生。
|
||||
|
||||
由于寄存器的高速读取特性，因此每次计算都离不开寄存器的参与。接着是一级缓存和共享内存，然后是常量内存、纹理内存、全局内存，最后则是主机端内存。根据不同存储器之间的存储速度的数量级的变化规律，选用适当类型的内存以及最大化地利用它们，从而发挥硬件的最大算力，减少计算时间。
|
||||
|
||||
### 硬件加速器的计算单元
|
||||
|
||||
为了支持不同的神经网络模型，加速器会提供以下几种计算单元，不同的网络层可以根据需要选择使用对应的计算单元。如图 :numref:`compute-unit`所示
|
||||
|
||||
- **标量计算单元**：与标准的精简指令运算集（Reduced Instruction Set Computer，RISC）相似，一次计算一个标量元素。
|
||||
|
||||
- **一维向量计算单元**：一次可以完成多个元素的计算，与传统的CPU和GPU架构中单指令多数据（SIMD）相似,已广泛应用于高性能计算（High Performance Computing，HPC）和信号处理中。
|
||||
|
||||
- **二维向量计算单元**：一次运算可以完成一个矩阵与向量的内积，或向量的外积。利用数据重复使用这一特性，降低数据通信成本与存储空间，更高效的提高矩阵乘法性能。
|
||||
|
||||
- **三维向量计算单元**：一次完成一个矩阵的乘法，专为神经网络应用设计的计算单元，更充分利用数据重复特性，隐藏数据通信带宽与数据计算的差距。
|
||||
|
||||
![多种计算单元[]](../img/ch06/compute_unit.svg)
|
||||
:width:`800px`
|
||||
:label:`compute-unit`
|
||||
|
||||
GPU计算单元主要由标量计算单元和三维向量计算单元组成。如图 :numref:`SM`所示,对于每个SM，其中64个32位浮点运算单元、64个32位整数运算单元、32个64位浮点运算单元均为标量计算单元。而8个张量核则是专为神经网络应用设计的三维向量计算单元。
|
||||
|
||||
![Volta GV100 流多处理器（SM） [@2017NVIDIA] []](../img/ch06/SM.svg)
|
||||
:width:`800px`
|
||||
:label:`SM`
|
||||
|
||||
张量核（Tensor Core）每个时钟周期完成一次4x4的矩阵乘累加计算,如图 :numref:`tensorcore`：
|
||||
|
||||
```
|
||||
D = A * B + C
|
||||
```
|
||||
|
||||
![Tensor Core 4x4矩阵乘累加计算[]](../img/ch06/tensor_core.svg)
|
||||
:width:`800px`
|
||||
:label:`tensorcore`
|
||||
|
||||
其中A,B,C和D都是4x4的矩阵，矩阵乘累加的输入矩阵A和B是FP16的矩阵，累加矩阵C和D可以是FP16也可以是FP32。 V100的张量核是可编程的矩阵乘法和累加计算单元，可以提供多达125 Tensor TFLOPS(Tera Floating-point Operations Per Second)的训练和推理应用。相比于普通的FP32计算单元可以提速10倍以上。
|
||||
|
||||
### DSA芯片架构
|
||||
|
||||
为了满足飞速发展的深度神经网络对芯片算力的需求，业界也纷纷推出了特定领域架构DSA芯片设计。以华为公司昇腾系列AI处理器为例，本质上是一个片上系统（System on Chip，SoC），主要应用在图像、视频、语音、文字处理相关的场景。主要的架构组成部件包括特制的计算单元、大容量的存储单元和相应的控制单元。该芯片由以下几个部分构成：芯片系统控制CPU（Control CPU），AI计算引擎（包括AI Core和AI CPU），多层级的片上系统缓存（Cache）或缓冲区（Buffer），数字视觉预处理模块（Digital Vision Pre-Processing，DVPP）等。
|
||||
|
||||
|
||||
:width:`800px`
|
||||
:label:`davinci_architecture`
|
||||
|
||||
昇腾AI芯片的计算核心主要由AI Core构成，负责执行标量、向量和张量相关的计算密集型算子。AI Core采用了达芬奇架构，基本结构如图 :numref:`davinci_architecture`所示，从控制上可以看成是一个相对简化的现代微处理器基本架构。它包括了三种基础计算单元：矩阵计算单元（Cube Unit）、向量计算单元（Vector Unit）和标量计算单元（Scalar Unit）。这三种计算单元分别对应了张量、向量和标量三种常见的计算模式，在实际的计算过程中各司其职，形成了三条独立的执行流水线，在系统软件的统一调度下互相配合达到优化计算效率的目的。 同GPU类似，在矩阵乘加速设计上，在AICore中也提供了矩阵计算单元作为昇腾AI芯片的核心计算模块，意图高效解决矩阵计算的瓶颈问题。矩阵计算单元提供强大的并行乘加计算能力，可以用一条指令完成两个16\*16矩阵的相乘运算，等同于在极短时间内进行了16\*16\*16=4096个乘加运算，并且可以实现FP16的运算精度。
|
||||
18
chapter_hardware_accelerator/accelerator_introduction.md
Normal file
18
chapter_hardware_accelerator/accelerator_introduction.md
Normal file
@@ -0,0 +1,18 @@
|
||||
概述
|
||||
----
|
||||
|
||||
### 硬件加速器设计的意义
|
||||
|
||||
未来人工智能发展的三大核心要素是数据、算法和算力。目前，人工智能系统算力大都构建在CPU+GPU之上，主体多是GPU。随着神经网络的层数越多，模型体量越大，算法越复杂，CPU和GPU很难再满足新型网络对于算力的需求。例如，2015年谷歌的AlphaGo与樊麾对弈时，用了1202个CPU和176个GPU，每盘棋需要消耗上千美元的电费，而与之对应的是樊麾的功耗仅为20瓦。
|
||||
|
||||
虽然GPU在面向向量、矩阵以及张量的计算上，引入许多新颖的优化设计，但由于GPU需要支持的计算类型复杂，芯片规模大、能耗高，人们开始将更多的精力转移到深度学习硬件加速器的设计上来。和传统CPU和GPU芯片相比，新型深度学习加速器会有更高的性能，以及更低的能耗。未来随着人们真正进入智能时代，智能应用的普及会越来越广泛，到那时每台服务器、每台智能手机、每个智能摄像头，都需要使用加速器。
|
||||
|
||||
### 硬件加速器设计的思路
|
||||
|
||||
近些年来，计算机体系结构的研究热点之一就是深度学习硬件加速器的设计。在体系结构的研究中，能效和通用性是两个重要的衡量指标。能效关注单位能耗下基本计算的次数，通用性主要指芯片能够覆盖的任务种类。
|
||||
|
||||
以两类特殊的芯片为例:一种是我们较为熟悉的通用处理器(如CPU)，该类芯片理论上可以完成各种计算任务，但是其能效较低大约只有0.1TOPS/W；另一种是专用集成电路(Application Specific Integrated Circuit, ASIC)，其能效更高，但是支持的任务相对而言就比较单一。对于通用的处理器而言，为了提升能效，在芯片设计上有许多加速技术的引入，例如：超标量技术、单指令多数据（Single Instruction Multiple Data，SIMD）技术以及单指令多线程（Single Instruction Multiple Threads，SIMT）技术等。
|
||||
|
||||
对于不同的加速器设计方向，业界也有不同的硬件实现。针对架构的通用性，NVIDIA持续在其GPU芯片上发力，先后推出了Volta, Turing, Ampere架构，并推出用于加速矩阵计算的张量核（Tensor Core），以满足深度学习海量算力的需求。
|
||||
|
||||
对于偏定制化的硬件架构，面向深度学习计算任务，业界提出了特定领域架构(Domain Specific Architecture)。 Google公司推出了TPU芯片，专门用于加速深度学习计算任务，其使用脉动阵列(Systolic Array)来优化矩阵乘法和卷积运算，可以充分地利用数据局部性，降低对内存的访问次数。华为也推出了自研的昇腾AI处理器，旨在为用户提供更高能效的算力和易用的开发、部署体验，其中的CUBE运算单元，就用于加速矩阵乘法的计算。
|
||||
162
chapter_hardware_accelerator/accelerator_programming.md
Normal file
162
chapter_hardware_accelerator/accelerator_programming.md
Normal file
@@ -0,0 +1,162 @@
|
||||
加速器基本编程原理
|
||||
------------------
|
||||
|
||||
本章前两节主要介绍了这些硬件加速器设计的意义、思路以及基本组成原理。软硬件协同优化作为构建高效AI系统的一个重要指导思想，需要软件算法/软件栈和硬件架构在神经网络应用中互相影响、紧密耦合。为了最大限度地发挥加速器的优势，要求能够基于硬件系统架构设计出一套较为匹配的指令或编程（操纵）方法。因此，本节将以[1.2.3](#硬件加速器的计算单元){reference-type="ref" reference="硬件加速器的计算单元"}中介绍的Tensor Core为例，着重介绍加速器的可编程性，以及如何通过编程使能加速器，提升神经网络算子的计算效率。
|
||||
|
||||
### 硬件加速器的可编程性
|
||||
|
||||
[1.1.2](#硬件加速器设计的思路){reference-type="ref" reference="硬件加速器设计的思路"}节中列出的硬件加速器均具有一定的可编程性，程序员可以通过软件编程，有效的使能上述加速器进行计算加速。但出于计算效率和易用性等方面考虑，将编程使能方式分为不同等级，一般包括：算子库层级，编程原语层级，以及指令层级。为了更具象的解释上述层级的区别，仍以Volta架构的Tensor Core加速器为例，由高层至底层对比介绍这三种不同编程方式：
|
||||
|
||||
- **算子库层级**：如cuBLAS基本矩阵与向量运算库，cuDNN深度学习加速库，均通过Host端调用算子库提供的核函数使能TensorCore；
|
||||
|
||||
- **编程原语层级**：如基于CUDA的WMMA API编程接口。同算子库相比，需要用户显式调用计算各流程，如矩阵存取至TensorCore、TensorCore执行矩阵乘累加运算、TensorCore累加矩阵数据初始化操作等；
|
||||
|
||||
- **指令层级**：如PTX ISA MMA指令集，提供更细粒度的mma指令，便于用户组成更多种形状的接口，通过CUDA Device端内联编程使能TensorCore。
|
||||
|
||||
矩阵乘法运算作为深度学习网络中占比最大的计算，对其进行优化是十分必要的。因此本节将统一以矩阵乘法$D[M, N] = C[M, N] + A[M, K] * B[K, N]$为实例，对比介绍如何通过不同编程方式使能加速器。
|
||||
|
||||
```
|
||||
$A, B$矩阵 $D$矩阵 $C[i][j] \gets 0$
|
||||
$C[i][j] \gets C[i][j] + A[i, k] \times B[k, j]$
|
||||
$D[i][j] \gets C[i][j]$
|
||||
```
|
||||
|
||||
### 硬件加速器的多样化编程方法
|
||||
|
||||
#### 算子库使能加速器
|
||||
|
||||
在上述三种层级的编程方式中，直接调用算子加速库使能加速器无疑是最快捷高效的方式。NVIDIA提供了cuBLAS/cuDNN两类算子计算库，cuBLAS提供了使能Tensor Core单元的接口，用以加速矩阵乘法(GEMM)运算，cuDNN提供了对应接口加速卷积(CONV)运算等。
|
||||
|
||||
以[1.3.1](#硬件加速器的可编程性){reference-type="ref" reference="硬件加速器的可编程性"}小节的GEMM运算为例，与常规CUDA调用cuBLAS算子库相似，通过cuBLAS加速库使能Tensor Core步骤包括：
|
||||
|
||||
1. 创建cuBLAS对象句柄且设置对应数学计算模式
|
||||
|
||||
```
|
||||
cublasHandle_t handle;
|
||||
cublasStatus_t cublasStat = cublasCreate(&handle);
|
||||
cublasStat = cublasSetMathMode(handle, CUBLAS_TENSOR_OP_MATH);
|
||||
```
|
||||
|
||||
2. 分配和初始化矩阵内存空间及内容元素
|
||||
|
||||
```
|
||||
size_t matrixSizeA = (size_t)M * K;
|
||||
cublasStat = cudaMalloc(&devPtrA[0], matrixSizeA * sizeof(devPtrA[0][0]));
|
||||
cublasStat = cublasSetMatrix(M, K, sizeof(A[0]), A, M, devPtrA[i], M);
|
||||
```
|
||||
|
||||
3. 调用对应计算函数接口
|
||||
|
||||
```
|
||||
cublasStat = cublasGemmEx(handle, transa, transb, m, n, k, alpha,
|
||||
A, CUDA_R_16F, lda,
|
||||
B, CUDA_R_16F, ldb,
|
||||
beta, C, CUDA_R_16F, ldc, CUDA_R_32F, algo);
|
||||
```
|
||||
|
||||
4. 传回结果数据
|
||||
|
||||
```
|
||||
cublasStat = cublasGetMatrix(M, N, sizeof(D[0]), devPtrD[i], M, D, M);
|
||||
```
|
||||
|
||||
5. 释放内存和对象句柄
|
||||
|
||||
```
|
||||
cudaFree(devPtrA);
|
||||
cudaDestroy(handle);
|
||||
```
|
||||
|
||||
当然，由于加速器一般会有矩阵形状、数据类型、排布方式等限制，因此在调用句柄和函数接口时要多加注意。如本例中，cuBLAS计算模式必须设置为$CUBLAS\_TENSOR\_OP\_MATH$，步长必须设置为8的倍数，输入数据类型必须为$CUDA\_R\_16F$等。按照如上方式即可通过cuBLAS算子库对[1.3.1](#硬件加速器的可编程性){reference-type="ref" reference="硬件加速器的可编程性"}实例使能Tensor Core加速器，通过NVIDIA官方数据可知，该方式对于不同矩阵乘法计算规模，平均有4～10倍的提升，且矩阵规模越大，加速器提升效果越明显。
|
||||
|
||||
该方式由于能够隐藏体系结构细节，易用性较好，且一般官方提供的算子库吞吐量较高。但与此同时，这种算子颗粒度的库也存在一些问题，如不足以应对复杂多变的网络模型导致的算子长尾问题（虽然常规形式算子占据绝大多数样本，但仍有源源不断的新增算子，因其出现机会较少，算子库未对其进行有效优化。），以及错失了较多神经网络框架优化（如算子融合）的机会。
|
||||
|
||||
#### 编程原语使能加速器
|
||||
|
||||
第二种加速器编程方式为编程原语使能加速器，如通过在Device端调用CUDA WMMA (Warp Matrix Multiply Accumulate) API接口。以线程束（即Warp，是调度的基本单位）为操纵对象，使能多个Tensor Core单元。该方式在CUDA 9.0中被公开，程序员可通过添加API头文件的引用和命名空间定义来使用上述API接口。基于软硬件协同设计的基本思想，该层级编程API的设计多与架构绑定，如WMMA操纵的总是$16x16$大小的矩阵块，并且操作一次跨两个TensorCore进行处理，本质是与TensorCore如何集成进SM中强相关的。针对Float16输入数据类型，NVIDIA官方提供了三种不同矩阵规模的WMMA乘累加计算接口，分别为$16x16x16$，$32x8x16$，$8x32x16$。
|
||||
|
||||
该API接口操纵的基本单位为Fragment，是一种指明了矩阵含义（乘法器/累加器）、矩阵形状（$WMMA\_M, WMMA\_N, WMMA\_K$）、数据类型（Half/ Float）、排布方式（$row\_major/ col\_major$）等信息的模板类型，包括如下：
|
||||
|
||||
```
|
||||
wmma::fragment<wmma::matrix_a, WMMA_M, WMMA_N, WMMA_K, half, wmma::row_major> a_frag;
|
||||
wmma::fragment<wmma::matrix_b, WMMA_M, WMMA_N, WMMA_K, half, wmma::col_major> b_frag;
|
||||
wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> acc_frag;
|
||||
wmma::fragment<wmma::accumulator, WMMA_M, WMMA_N, WMMA_K, float> c_frag;
|
||||
```
|
||||
|
||||
使用时，我们需要将待执行乘法操作矩阵块的数据，作为Fragment，由寄存器加载至TensorCore，在将累加Fragment初始化/清零操作后，通过TensorCore单元执行乘累加运算，最后将运算结果的Fragment存回寄存器或其他内存区域。与上述操作对应的，NVIDIA提供了$wmma.load\_matrix\_sync(), wmma.store\_matrix\_sync()$接口用于将参与计算的子矩阵块写入/载出Fragment片段；$wmma.fill\_fragment()$接口用于初始化对应Fragment的数据；$wmma.mma\_sync()$接口用于对Fragment进行乘累加运算。
|
||||
|
||||
#### 指令集编程使能加速器
|
||||
|
||||
在NVIDIA PTX ISA (Instruction Set Architecture)中提供了另一个编程接口，如Volta架构中的$mma.sync.m8n8k4$指令，它使用$M=8, N=8, K=4$的形状配置执行乘累加操作。具体地，它由线程组（黑色椭圆表示）或octet执行[@2018Modeling]，如图 :numref:`PTX`显示了线程和数据的映射关系。每个线程组由四个连续的线程组成，使用不同颜色的圆圈表示。图中还指出了一个octet里面的线程在线程束内的分布，Float16乘法器A或B的四个连续元素（使用具有相同颜色的块表示），以及Float32累加器C或D的八个分散元素（同样使用相同颜色的块表示）。彩色块上的数字代表对应的线程ID。
|
||||
|
||||
|
||||
:width:`800px`
|
||||
:label:`PTX`
|
||||
|
||||
作为一个更细粒度的指令，mma可以组成更加多样化形状的Warp范围的WMMA API接口，可以控制线程束内线程与数据的映射关系，并允许AI编译器自动/手动显式地管理内存层次结构之间的矩阵分解，因此相比于直接应用NVCUDA::WMMA API具有更好的灵活性。
|
||||
|
||||
#### 算子编译器编程使能加速器
|
||||
|
||||
除了上述三种层级的编程方式，算子编译器也逐渐成为DSA加速器的关注热点。随着深度学习模型的迭代更新以及各类DSA加速器的层出不穷，手写算子或高性能算子库（如cuDNN/cuBLAS）等基于人工优化算子的方式给算子开发团队带来沉重的负担。因此，开发一种能够将High-level的算子表示编译成目标硬件可执行代码的算子编译器，成为了学术界、业界的共识。
|
||||
|
||||
近年来涌现出许多算子编译器/编译框架，如TVM，为不同的硬件后端提供了编译优化支持。在昇腾芯片上，基于TVM开发了TBE(Tensor Boost Engine)，不仅提供了一个优化过的神经网络标准算子库，同时还提供了算子开发能力及融合能力。通过TBE提供的API和自定义算子编程开发界面可以完成相应神经网络算子的开发，帮助用户较容易的去使能硬件加速器上的AI\_CORE相关指令，以实现高性能的神经网络计算。此外，为了更好的支持复杂算子融合场景，还提供了基于polyhedral多面体编译技术的AKG(Auto kernel generator)，提供算子的自动生成能力。
|
||||
|
||||
基于算子编译器使能加速器实现矩阵乘的流程则对用户更加友好，用户只需基于python定义矩阵乘的tensor信息（数据类型及形状等），调用对应TBE接口即可。如下所示：
|
||||
|
||||
```python
|
||||
a_shape = (1024, 256)
|
||||
b_shape = (256, 512)
|
||||
bias_shape = (512, )
|
||||
in_dtype = "float16"
|
||||
dst_dtype = "float32"
|
||||
tensor_a = tvm.placeholder(a_shape, name='tensor_a', dtype=in_dtype)
|
||||
tensor_b = tvm.placeholder(b_shape, name='tensor_b', dtype=in_dtype)
|
||||
tensor_bias = tvm.placeholder(bias_shape, name='tensor_bias', dtype=dst_dtype)
|
||||
res = te.lang.cce.matmul(tensor_a, tensor_b, False, False, False, dst_dtype=dst_dtype, tensor_bias=tensor_bias)
|
||||
```
|
||||
|
||||
### 硬件加速器高性能编程实例
|
||||
|
||||
本节[1.3](#加速器基本编程原理){reference-type="ref" reference="加速器基本编程原理"}前几个小节主要介绍了硬件加速器的不同层级的多样化编程方法。调用计算库的方式留给程序员的优化空间较少，合理利用硬件加速器不同层级的编程，可以实现更好的性能优化。 为了更好的让读者理解硬件加速器的使用，本节会继续[1.3.1](#硬件加速器的可编程性){reference-type="ref" reference="硬件加速器的可编程性"}节中的GEMM运算，仍以WMMA API使能Tensor Core加速单元为例，介绍如何通过矩阵分块、资源映射等方式更高效的利用硬件加速器。
|
||||
|
||||
[\[alg:TensorCore\]]{#alg:TensorCore label="alg:TensorCore"}
|
||||
|
||||
```
|
||||
$A, B$矩阵 $D$矩阵 Mapping to Block.Idx Mapping to Block.Idy Mapping
|
||||
to Block.Idz
|
||||
$A_{Shared}[i_o][k_o] \gets A[i_o][k_o]$ $B_{Shared}[k_o][j_o] \gets B[k_o][j_o]$ $Syncthreads()$
|
||||
Mapping to Warp.Idx Mapping to Warp.Idy
|
||||
$A_{Register}[i_i][k_i] \gets A_{Shared}[i_i][k_i]$ $B_{Register}[k_i][j_i] \gets B_{Shared}[k_i][j_i]$ $pragma\ unroll$
|
||||
$wmma.load\_matrix\_sync(A_{Fragment}, A_{Register})$ $wmma.load\_matrix\_sync(B_{Fragment}, B_{Register})$
|
||||
$wmma.fill\_fragment(C_{Fragment}, 0)$
|
||||
$wmma.mma\_sync(D_{Fragment}, C_{Fragment}, A_{Fragment}, B_{Fragment})$
|
||||
$Syncthreads()$ $wmma.store\_matrix\_sync(D, D_{Fragment})$
|
||||
```
|
||||
|
||||
若要得到高性能CUDA程序，提高并行性、增大吞吐量、优化指令执行是至关重要的三个优化目标。针对该实例，具体地实现和优化方案列出如下，对应到具体实例伪代码如算法2所示：
|
||||
|
||||
1. **优化内存结构------增大吞吐量**：将原始大规模矩阵根据不同阈值切分成不同层级的子矩阵块，使得子矩阵块能被如共享内存、寄存器等高性能体系结构存储下来，以此提高吞吐量。设置切分参数为$BlockTile[Ms, Ns, Ks]$和$WarpTile[Mw, Nw, Kw]$，对应的将BlockTile下的矩阵由全局内存搬移至共享内存，以提高全局内存合并访问和数据局部性，如图 :numref:`GEMM-BlockTile`所示；再将内层WarpTile下的矩阵由共享内存搬移至寄存器中，如图 :numref:`GEMM-WarpTile`所示，以备Tensor Core加速器数据存取。
|
||||
|
||||
|
||||
:width:`800px`
|
||||
:label:`GEMM-BlockTile`
|
||||
|
||||
|
||||
:width:`800px`
|
||||
:label:`GEMM-WarpTile`
|
||||
|
||||
2. **并行资源映射------提高并行性**：将多层级的并行资源（Block、Warp、Thread）与对应需要计算/搬移的数据建立映射关系，提高程序并行性。将可并行的计算/数据搬移操作映射到并行资源上，对于GEMM实例，M/N轴即为可并行轴，将数据搬移操作中的循环指令映射分配到Block层级（即算法3中的2-4行$For$循环），将内层循环指令映射分配到Warp层级（即算法3中的8-10行$For$循环）。（前文介绍，线程束Warp作为调度的基本单位，且是WMMA API操纵的基本层级，因此对Warp层级进行数据映射比Thread层级映射更为合适）
|
||||
|
||||
3. **Warp统一的Tensor Core数据交互------增大吞吐量**：根据[1.3.2](#硬件加速器的多样化编程方法){reference-type="ref" reference="硬件加速器的多样化编程方法"}节中介绍的编程方法，除调用算子库外，均需要使用或将指令封装成WMMA接口形式统一进行Warp层级的数据存取和计算。如图 :numref:`GEMM-TensorCore`所示，Tensor Core加速器需要从局部内存/寄存器中读取数据，存于虚拟Fragment数据结构中，对应使用$wmma.load\_matrix\_sync()$接口，将累加Fragment $C$ 通过$wmma.fill\_fragment()$接口进行初始化后，使用$wmma.mma\_sync()$使能加速器进行乘累加运算，后将结果Fragment $D$通过调用$wmma.store\_matrix\_sync()$接口拷贝至目标内存地址。
|
||||
|
||||
|
||||
:width:`800px`
|
||||
:label:`GEMM-TensorCore`
|
||||
|
||||
4. **优化数据访存------提高并行性**：在进行内存结构变化（矩阵数据搬移）时，需要注意全局内存的合并访问、共享内存的存储体冲突等常见性能瓶颈点。
|
||||
|
||||
5. **资源负载均衡------增大吞吐量**：调整平衡每个线程处理的数据量、共享内存使用量、寄存器使用量，以获得更高的SM占用率。一般在实际程序中BlockTile和WarpTile的选取至关重要。
|
||||
|
||||
6. **优化指令执行**：使用\#unroll功能进行循环展开以避免分支冲突，如算法3中13行；使用向量化加载指令减少PTX指令执行次数以提高带宽等，对于GPU Volta架构，最大向量化加载指令为ldg128，即128比特带宽，对于算法3中5-6行数据由全局内存加载至共享内存时，即可采用Float4\*类型指针进行内存读取。
|
||||
|
||||
@@ -1,10 +1,24 @@
|
||||
# 硬件加速器
|
||||
|
||||
上一章节，我们详细讨论了计算图的基本组成，生成和执行等关键设计。当前主流深度学习模型大多基于神经网络实现，无论是训练还是推理，都会产生海量的计算任务，尤其是涉及矩阵乘法这种高计算任务的算子。另一方面，通用处理器芯片如CPU，在执行这类算子时通常耗时较大，难以满足训练/推理任务的需求。因此工业界和学术界都将目光投向特定领域的加速器芯片设计，希望以此来解决算力资源不足的问题。
|
||||
|
||||
本章将会着重介绍加速器的基本组成原理，并且以矩阵乘法为例，介绍在加速器上的编程方式及优化方法。最后，介绍由异构算子组成的异构计算图表达与执行方式。
|
||||
|
||||
本章的学习目标包括：
|
||||
|
||||
- 掌握加速器的基本组成
|
||||
|
||||
- 掌握矩阵乘法的常见优化手段
|
||||
|
||||
- 理解编程API的设计理念
|
||||
|
||||
- 理解异构硬件加速的表达与执行
|
||||
|
||||
```toc
|
||||
:maxdepth: 2
|
||||
:numbered:
|
||||
|
||||
|
||||
```
|
||||
accelerator_introduction
|
||||
accelerator_architecture
|
||||
accelerator_programming
|
||||
summary
|
||||
```
|
||||
|
||||
10
chapter_hardware_accelerator/summary.md
Normal file
10
chapter_hardware_accelerator/summary.md
Normal file
@@ -0,0 +1,10 @@
|
||||
总结
|
||||
----
|
||||
|
||||
- 面向深度学习计算任务，加速器通常都是由多种片上缓存以及多种运算单元组成来提升性能。
|
||||
|
||||
- 未来性能增长需要依赖架构上的改变，即需要利用可编程的硬件加速器来实现性能突破。
|
||||
|
||||
- 出于计算效率和易用性等原因，加速器一般会具有多个等级的编程方式，包括：算子库层级，编程原语层级和指令层级。
|
||||
|
||||
- 越底层的编程方式越能够灵活地制加速器，但同时对程序员的能力要求也越高。
|
||||
1233
img/ch06/G2S.svg
Normal file
1233
img/ch06/G2S.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 90 KiB |
2880
img/ch06/R2TC.svg
Normal file
2880
img/ch06/R2TC.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 214 KiB |
1193
img/ch06/S2R.svg
Normal file
1193
img/ch06/S2R.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 87 KiB |
9571
img/ch06/SM.svg
Normal file
9571
img/ch06/SM.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 717 KiB |
22683
img/ch06/V100.svg
Normal file
22683
img/ch06/V100.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 1.7 MiB |
13867
img/ch06/compute_unit.svg
Normal file
13867
img/ch06/compute_unit.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 1.0 MiB |
3715
img/ch06/davinci_architecture.svg
Normal file
3715
img/ch06/davinci_architecture.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 277 KiB |
1265
img/ch06/ptx.svg
Normal file
1265
img/ch06/ptx.svg
Normal file
File diff suppressed because it is too large
Load Diff
|
After Width: | Height: | Size: 92 KiB |
392
img/ch06/tensor_core.svg
Normal file
392
img/ch06/tensor_core.svg
Normal file
@@ -0,0 +1,392 @@
|
||||
<?xml version="1.0" encoding="UTF-8" standalone="no"?>
|
||||
<!-- Created with Inkscape (http://www.inkscape.org/) -->
|
||||
|
||||
<svg
|
||||
version="1.1"
|
||||
id="svg55"
|
||||
width="1024"
|
||||
height="254"
|
||||
viewBox="0 0 1024 254"
|
||||
sodipodi:docname="tensor_core.svg"
|
||||
inkscape:version="1.1.2 (1:1.1+202202050950+0a00cf5339)"
|
||||
xmlns:inkscape="http://www.inkscape.org/namespaces/inkscape"
|
||||
xmlns:sodipodi="http://sodipodi.sourceforge.net/DTD/sodipodi-0.dtd"
|
||||
xmlns:xlink="http://www.w3.org/1999/xlink"
|
||||
xmlns="http://www.w3.org/2000/svg"
|
||||
xmlns:svg="http://www.w3.org/2000/svg">
|
||||
<defs
|
||||
id="defs59" />
|
||||
<sodipodi:namedview
|
||||
id="namedview57"
|
||||
pagecolor="#ffffff"
|
||||
bordercolor="#666666"
|
||||
borderopacity="1.0"
|
||||
inkscape:pageshadow="2"
|
||||
inkscape:pageopacity="0.0"
|
||||
inkscape:pagecheckerboard="0"
|
||||
showgrid="false"
|
||||
fit-margin-top="0"
|
||||
fit-margin-left="0"
|
||||
fit-margin-right="0"
|
||||
fit-margin-bottom="0"
|
||||
inkscape:zoom="1.5966797"
|
||||
inkscape:cx="296.55291"
|
||||
inkscape:cy="166.90887"
|
||||
inkscape:window-width="1848"
|
||||
inkscape:window-height="1016"
|
||||
inkscape:window-x="72"
|
||||
inkscape:window-y="27"
|
||||
inkscape:window-maximized="1"
|
||||
inkscape:current-layer="g61" />
|
||||
<g
|
||||
inkscape:groupmode="layer"
|
||||
inkscape:label="Image"
|
||||
id="g61"
|
||||
transform="translate(0,1282.8716)">
|
||||
<image
|
||||
width="1024"
|
||||
height="254"
|
||||
preserveAspectRatio="none"
|
||||
xlink:href="data:image/png;base64,iVBORw0KGgoAAAANSUhEUgAABAAAAAD+CAYAAACkynf5AAAABHNCSVQICAgIfAhkiAAAIABJREFU
|
||||
eJzt3X9sled9//+XvwpRVVs79kfgAyTCuLZJDNjG/NF0hwGyq+DYjMjLJlqqEqNsM5m6ymYJS/+g
|
||||
paSRFsYSbFVRsaUyTKumRVtqwWICUUEkw1oqreZnCNiuMUoCx6DaluytCpPO94/zw+ccn2Of43Pu
|
||||
n+f5kCz5HN/n+Hr59nmfc1/3dV+XAgDmNTU1FfjZz34WOHLkyJzbnT59OrBy5cqAJNd8PfLII5a3
|
||||
gTy5k2flypWBf/3Xf034+nrvvfcCP/3pTwPj4+MGvMozR52wvh3kyY08Tq4TbpRq7QsEqH9O+CKP
|
||||
vb+S1b+TJ08Gjhw5EvD7/fO+DvMCgUBAAJL67LPP9P3vf1/j4+PKz8/Xz3/+86TbLlu2TPfu3TOx
|
||||
dYD7fOlLX9L//u//zrr/+9//vq5du6b8/Hy99NJLqq2ttaB1iVEnAHM5sU64UTq1T6L+AdmQqP59
|
||||
73vf082bN1Oqff+f0Q0EnGx6elqvvvqqxsfHI7dHRkaSbs+bGpC5P/7xj7Pue/jwoW7duiUp+Dp8
|
||||
/fXXde3aNbOblhB1AjCf0+qEG6Vb+yTqH5ANiepf+LU3PT2tN954Q5999lnSxz9iWMsAF/j5z3+u
|
||||
sbGxmPtGRkZUWlo6/4OP/dCYRplp1w9nvieP/bg5T5ybN2/qiy++iNz+4osv9Oabb+rNN99UYWGh
|
||||
CY1LLpM60b7umEGtMk/HpV2R78ljP27OE8/OdcKNMvqMJOkfzxjRKnP9c8PM9+SxHzfniTYyMhJT
|
||||
+8KdAG+++WbC7RkBACQxMjKi9957L+a++vp6rV271qIWAbnr8ccf186dO2PuGx8f169+9SuLWhRE
|
||||
nQDsw651wo2ofYB9FBUV6S//8i9j7hsZGdFHH32UcHs6AIAkzp8/H3P78ccf14svvqji4mKLWgTk
|
||||
rsLCQj333HOzPty/9957s85AmYk6AdiHXeuEG1H7APsoLCzUt7/9bT311FMx9//6179OuD0dAEAC
|
||||
Dx8+1H/+53/G3Pftb39bixYtsqhFACRp27ZtKioqirnvP/7jPyxpC3UCsCc71Qk3ovYB9vTCCy/E
|
||||
3L5582bCeTnoAAAS+OijjyKT2kjBoTXxvWoAzLdo0SL9xV/8Rcx9586d08OHD01vC3UCsCc71Qk3
|
||||
ovYB9lRcXDzrtRg/WkeiAwBI6MyZ2JlCtmzZYlFLAMR75plnlJ+fH7k9PT1tyUzf1AnAvuxSJ9yI
|
||||
2gfYV11dXcztRJ2fdAAAcaanpyPLCIU988wzFrUGQLxFixZp48aNMffFD0c1GnUCsDc71Ak3ovYB
|
||||
9vbUU0/FXAI1PT2tmzdvxmxDBwAQ59q1azFLaZSWlrJ8EGAzGzZsiLlt9pk96gRgf1bXCTei9gH2
|
||||
V1tbG3P7t7/9bcxtOgCAONevX4+5/dWvftWilgBI5oknntCjjz4auT02NmbqLN/UCcD+rK4TbkTt
|
||||
A+wv/nUZ3/lJBwAQJ37NTNa0Bexn0aJFWrVqVcx9Zp7do04A9md1nXAjah9gf/Gvy5GREU1PT0du
|
||||
0wEARIk/O/Doo4/qiSeesLBFAJKJH+IWf2bKKNQJwDmsqhNuRO0DnCE/P1+lpaUx90V3ftIBAESJ
|
||||
n9hm7dq1rGsL2FT8B/v4M1NGoU4AzmFVnXAjah/gHPGjAAYHByPf0wEARLl9+3bM7fjeMwD2UVpa
|
||||
OmuZr88++8zw30udAJzDqjrhRtQ+wDniL38aGRmJfE8HABDl008/jbm9cuVKaxoCICXxw0+j3+CM
|
||||
Qp0AnMWKOuFG1D7AOeI76OgAAJKIPyvw2GOPWdQSAKmIf4OLP0NlBOoE4CxW1Ak3ovYBzvHYY4/F
|
||||
rIIyPj4emQiQDgAg5OHDhzG9248++qgef/xxC1sEYD7xZ6Diz1BlG3UCcB6z64QbUfsA54nv/AzP
|
||||
40EHABAS/4HgscceY3IbwObmGuJmBOoE4Dxm1wk3ovYBzhNf+8Kv40esaAxgR/FD26zo2W6t268u
|
||||
z4fa0HtO/ab/9uwjj725IU94iNsXX3whKbhM1fT0dMykX9lkZZ1Y39GinZsT/GD0ig41D+hz01qS
|
||||
HeSxNzflMbtOuJEdPiPJ26eDNY0xd12/nKfjfvObkhXksTcX5Ikf/fT558HKzQgAICR6bVtJWr58
|
||||
ucktqFLzCkmeSm0vMvlXG4I89uaePPE93EbO8G19nUigpFp7B+q03up2ZAt57M2hecysE25kbe3z
|
||||
qcEXmHUwJklragI6WNNqYluygTz25p48xcXFMbfDr2M6AICQ+/fvx9yOf9EYzVe7ScFSs1hbV3pN
|
||||
/d1GII+9uSmPmcN7ra4TkqQLH2hPbU/w6wd3Qneu0JY2h/bkkMfeXJKHywAyY13t86nBd1H1BcFb
|
||||
Y8Mb9MrZvODX5dPBO/NXa5lJrckceezNXXniX6fj4+OS6AAAIuJ7t839YO/V9pLFkVvlJWvkM/G3
|
||||
Zx957M1deZYsWRJzO/6DajZZWycSOHVbV61tQXaRx94cnMfMOuFGltU+777Iwdj1y3l6YzjqgjV/
|
||||
kzqGhzTmP6G75rQmc+SxN5flYQQAMA9LP9gXrdFWj6TJGzo9KcmzUftK53uQjZHH3lyWJ1kPtxFs
|
||||
1wGwbaWqJEmTunLGuNymIY+9OTiPmXXCjaypfT41lIWGYU916v0E117f9R/Sb/xOmcWGPPbmtjzS
|
||||
okWLYuY6mZ6e1vT0NJMAAmFWfrBvXbdR5ZKGRi/oNS1RY/ViNa6skkacea6FPPbmtjzJeriNYIsO
|
||||
gM2bdHhg08xtB07KFoM89uaSPGbWCTeypvatlTc8FDvZWdepbl0yoSXZQR57c1ueoOLi4phLniYm
|
||||
JhgBAEjBF0N4dmDJ7A/1ocnY9EDv3var//YNDUnSiko5Z5qRaOSxN7flMe+DvbV1Yg4l1do70KIX
|
||||
tlndkCwhj705NA8dAAtnWe0rWK3whRv3p5xz1jUp8tib2/KEJKp9dAAAsvisXmllcDK2yRs6MS5p
|
||||
/LrenZSkSu2tdeDkbOSxN7flkVRYWKhHH300ctuoD/a2OPsvxU7KVvtB5JrsqledNzO7JPLYnUvy
|
||||
mFUn3Mg2tQ9A2oqKYidsHR8fpwMAkBTTsy3JxLWBveqoqZQkDY1eD63F7teJ0QeSnDg5G3nszW15
|
||||
ZhQWFsbcNuLDvXV1Yi4jOhqZmd2jpU9a2pgsII+9OTuPGXXCjexQ+5YUOPXdKTHy2Jub8ng8npjb
|
||||
ExMTzAEASLM/BJj25haejE1SefWLClTH/dxTqe1F59QfnquodLsCmyojPx66ckQVAwlmKUl322wx
|
||||
Mk9Y6XYFah5oQ+85GT5Ay8A8vtrv6GJ1eCb+B+o8+ZbajZ6Tyqw8kx+as3+iFBcXx7yOx8bGsn6W
|
||||
yrI6MZ/BCfm1Ql55tLRC0ieSttXp8KsrIpv4j53U651J/sHS2dYMmeYJ21anw389Yf316hnmWd7W
|
||||
rL27wh/gJvX+jl71fWJ4q5PLZh6T5xMwo064kWW1b+qErk21qb5AKi5YK6XyruLti1m/fWx4Q+xM
|
||||
7gvdNhuMzhP9mLJb6uhvN3a2eoPzLCsbVHtZeejWkM71V+jMVObNTsrMPFOdxu+fkOiRT1KwQ48R
|
||||
AICkhw8fxtw2683Nt7JS5XNusVht66pC31epb1Olhq4cUV7PAW248kDl1X+ljoRLMaezbfYYlyeo
|
||||
tW5/zAGp0YzcP/uqpc6TB5TXc0C77yxW2+Z6w8++G5fHq+0l4Tzv6LRnoy7WVSXa0DDxH+InJiay
|
||||
/jusqhPzqihU8OKNSd0blKRSvfDqCvmPndSe2h4dOjYp767Nakp4tjadbU2SUZ6g9R0tMQeklspw
|
||||
/zyzS3p/R3AI/s8uePT067Vabl7rZ8soT5HW1YXzfKCrJdXa22HeEiRm1Ak3sq729euKfyj4rXev
|
||||
GgoSbdOqhjJf5Pvnaxoja7d3DA+puKwn6eNS3zZbjMwTtK4mEHNAaixj98/TZdK5/jy9cjZPb/vL
|
||||
VV/doWWG5AgzMo9P1d5wnt26XtCm9hpzZl6KH/k0OTlJBwAgzV4O6Mtf/rIJv7VK+0JnTMMHWdFf
|
||||
G64Eh2VHJmcrrVRjaOI2SaHJ2xZrVWGCp05nWyfkUfAMc/PtA9p9J/HPs8/IPFfV1DNzxv/axAPJ
|
||||
s1hrHZvHr/becJ4x3ZqU5Fli6uUE8R9I44esZoM1dWI+wYMvSdLoqC59otBSbTPLtH1+ZlT+8Nna
|
||||
eOlsa4oM8yh4hnndb3r0swvmtHhumeYZ0dHamTP+90YmpZJCLTWj6Qllmmdcfc3hPBO6NyppZaFp
|
||||
HRpm1Ak3srL23R0+pOuSpHLV+wJ6PmaqmlY9v6VL9eGDLm+z1mhI10LLtN31v6sxlcubqL8inW2z
|
||||
yLA8Cp5hrr6Xp7cNHuwZzbg83Tp+duaMv39qSCpYJaNnKjIuT7/OREYwXJN/SlL+aoM7NBL74osv
|
||||
uAQAkKT/+Z//ibltSu92eDK2qIOsaP23b2ioeqPKVanmUula4RJJ9/Vx+H14/L4GJVUUeiXFPt6X
|
||||
xrZZY2AeSeofeEtNklpXGtP8WQzOMyN0YD55Q9eyHCGGaXmKtcoj6c59Uy8BiP9AasS1vZbUiUTi
|
||||
l2WTJE3q/e8Fh1Iv/4pH0qTuhYeJfzKhMUnFXymSFPtBPp1tDZPFPJL0eWevjkpa/3UjGz2HLOeZ
|
||||
Uapndnmk0VHdM6DZSRmWp1BLSyRdmDDtEgAz6oQbWVv7unW8f7Ve8rWpWNKamoAOJtlyWUGFpMHg
|
||||
wZUkTX2s+wpdzx23dns622aXMXkk6e5whY5LWmdqD6FxeWa06umycmnqXaM+wUYxI09ouUH/x6Zc
|
||||
ApBo5BMjAABJ09PTMbeNf3ObmYwtMht7vMjs7FJjTb0sOyGXEvIsNI+vdpMaFT0pnxFMylO6XYGW
|
||||
59R45x3lnb86//ZZFH+NmxHMrxMpGr2iQ7UWXxeeTeRJaHlbjaok+c/ftnZOg2zk2VanwwObVHXh
|
||||
A+1pH5l/+ywxo064keW1b6pdb5xNfGY7OPza4GvDs408aVlWtldrJI35T5hywGxoHm+fDm7p0hr/
|
||||
br1yuTujZi4UIwCAkPjrAOOXzMg+v9p7D6g9jW18teZd+54+8iwoT1G9eqoXS5MfqsXQCRpNyjNy
|
||||
QnkjXnU0v6hA8xJTJwKMv8bt/v37Wf8d5teJWL9r79HvTP2NxiJPGp6s1fO7PNLoFR03aYJGQ/Oc
|
||||
Oq89p4rU1PusDvcWmjYRoBl1wo2srn1hly7n6ZIlv9kY5ElBQYe+VVYuTXXqF0ZO0JiAIXn8TXrl
|
||||
rE8Nvos66FttykSA8a9XRgAAIfHXAS5atMiiliTXP3Ff0hKtDr+Oi5aoQtLgRILh3GlsaxUntDEd
|
||||
6efxqmPzRpVLOn3Z3BnzU7Hw/RNaVtBTqe0mfkY048yeE+qEJH3++0nFLNH2ZKGKJY39PsFw+TS2
|
||||
tYoT2piO9PMUqen1anklXf2pxSsaJLDw/TOuS+cnpZISrTNp0klGACyMU2rf3alBSRXB4dWSVLBa
|
||||
SyTdn0owXD6Nba3ihDamI/08PjVUB4fiXx82Z8b8dCx8/4QmGyzYqmpDJ50Min+9sgoA4CQTDzSk
|
||||
xdq6MjgjSXBG9xvqHZGkKvW17FegOTST/Jzb2kQ6eZwg3Tylm9XmkXTnHTXZab+EpZXHq47m7QrP
|
||||
Z7u2cLFi5g8wAet7RxmckF8eVTcEe2CWN5TIqzu6dEqSSvXCQIsO94Zmkp9zW5tIJ48TpJtn2zo9
|
||||
XSLpwgc6aqf9EpZWniI19dZpfeihS0vj5g8wGHXC5aZvaUzlWusNvjMt825VsU7ril8KTuAW0EFf
|
||||
aCb5Obe1iXTyOEG6ebz7VF8gyb9bx+20X8LSyuNTg69P60IP9RaUK2b+AJNxCQDgFOPn1HKlUhej
|
||||
1m8//cFbCl5BdFW9d55T44rgTPL9c25rE+nksa6VqUsrT3CJPUnSiucUaHlOUnB2/gpDLwVIQ9r7
|
||||
p1JdLfvVJUl6oM6TJ+z1/5ZLPhnQ8WMl2rvrWR3eFbzr6g96Q8O5R3TpwiZVbQ7OJP/5nNvaRDp5
|
||||
LGtkGtLKEzXrftRkfP5jJ/W6SZcCzCvt/bNCOwdatFOSNKn3d5y31/8bnGuqXb8Y3qr2sos6WBa8
|
||||
6/rlitAw7m5d8XdpjTc4k/zdObe1iXTyWNfK1KWVJ7jEniTJ26WDW4KfLsaGN+gNky8FSCrt/dOo
|
||||
HVsC2iFJGtK5/ibL/t/oAAA0+yxA/IyZdtE/8JbyBubY4M6NyEHXvNvaQDp5wrrPH7DtgWXqea6q
|
||||
qcfcSfIWIvU8qcwxYKz416wRZ/acUiek4Ez4ezrn2ODC7chB17zb2kA6ecLsPMdA6nlGdLTWjkOE
|
||||
YqWeZ1x9zT3qM6dZs5hRJ9zISbXv7nCFXhmeYwN/b+Sga95tbSCdPGF2nlsg9TzdOn7Wrp/2ZqSe
|
||||
p19n+vN0xpxmxUhU97gEAHCD0u3qkvkzrxuGPPbmtjy5ZFuddsrcmdcNRR57c1seOJe3Tztk3czr
|
||||
WUcee7N5HkYAAG4wckJ5bvp8RR57c1ueXHLqvPbY8TryhSKPvbktD5zL36RXbHKFXVaQx95snocR
|
||||
AAAAAAAA5AA6AAAAjse1vQDmQ50AADoAAEnS9PR0zO38/HyLWgLArqgTAHIRtQ9wFzoAAPHmBmB+
|
||||
1AkAuYjaB7gLHQAAAAAAAOQAOgAAAAAAAMgBdAAAAAAAAJAD6AAAAAAAACAH0AEAAAAAAEAOoAMA
|
||||
AAAAAIAcQAcAAAAAAAA54BGrGwC41q4fWt2C7CKPvbktT47ouLTL6iZkFXnszW154Gz/3GB1C7KL
|
||||
PPbmtjyZYAQAkEV5eXlWNwGAzVEnAOQq6h9gPToAgCwKBAJWNwGAzVEnAOQq6h9gPS4BAIxy7IdW
|
||||
tyBz0cPKyWM/bs6TI9rXHbO6CRmLHlZOHvtxcx442z+esboFmYseVk4e+3FznkwwAgAAAAAAgBxA
|
||||
BwAAAAAAADmADgAAAAAAAHIAHQAAAAAAAOQAOgAAAAAAAMgBdAAAAAAAAJAD6AAAAAAAACAH0AEA
|
||||
AAAAAEAOoAMAAAAAAIAcQAcAAAAAAAA5gA4AAAAAAAByAB0AAAAAAADkADoAAAAAAADIAXQAAAAA
|
||||
AACQA+gAAAAAAAAgB9ABAAAAAABADnjE6gYAmNFat19dng+1ofec+q1uTBaQx97clsft1ne0aOfm
|
||||
BD8YvaJDzQP63PQWZYY89ua2PHABb58O1jTG3HX9cp6O+y1qT6bIY29uyxOFEQCAbVSpeYUkT6W2
|
||||
F1ndlmwgj725LU8OK6nW3oE6rbe6HdlCHntzWx44gE8NvsCsgzFJWlMT0MGaVgvalAny2Jvb8sxG
|
||||
BwBgE77aTQqWmsXautJrcWsyRx57c1uenHLhA+2p7Ql+/eBO6M4V2tLm0J4c8tib2/LAYXxq8F1U
|
||||
fUHw1tjwBr1yNi/4dfl08M781VpmXQPTRB57c1uexOgAAGzBq+0liyO3ykvWyGdhazJHHntzW54c
|
||||
duq2rlrdhmwij725LQ/sz7svcjB2/XKe3hiOumDN36SO4SGN+U/orjWtSx957M1teZKgAwCwg6I1
|
||||
2uqRNHlDpycleTZqX6nVjcoAeezNbXly2baVqpIkTerKmXGLG5MF5LE3t+WBzfnUUBYahj3VqfcT
|
||||
XHt9139Iv/E7ZRYb8tib2/IkxySAgA20rtuocklDoxf0mpaosXqxGldWSSPOPNdCHntzW56cs3mT
|
||||
Dg9smrnt9EnZyGNvbssDB1krb3godrKzrlPdumRmkzJCHntzW57k6ACQ9MYbbyS8f2hoSEeOHJnz
|
||||
sf/yL/8iSaqtrdXjjz8ur9crj8eT9TbCzUKTsemB3r3tV79uaKh6o8pXVKpVV9VtdfPSRh57c1se
|
||||
BCdlq9bVH/To6CmrG5MF5LE3t+WBfRWs1pLQt/ennH/WlTw257Y8c6ADQNLLL7+c1cfu27dPW7du
|
||||
1de+9rVMmoVcUVoZnIxt8oZOjEvSdb07uVFtnkrtrfWqe8Bh642Qx97clicXXfhAe9pHQjdK9cLA
|
||||
JlVJqnq1TutPndfvrGzbQpDH3tyWxyZ++9vf6u23317w4xsaGvTMM89ksUUAcgVzABjgtdde05/+
|
||||
6Z/q2Wef1blz56xuDmzNq46aSknS0Oj10Frsfp0YfSDJiZOzkcfe3JYH0oiORmZm92jpk5Y2JgvI
|
||||
Y29uy2OdP/zhD7p06dKCv+7du2d1BFMtKXDXuxN57M1teeLRAWCgU6dO6etf/7qeffZZ3bp1y+rm
|
||||
wI7Ck7FJKq9+UYGW/Qq07NfF6tAM7cnWaC/drkBz/dwHa6XbI88XaNmvwVoTlnozMs9Cts2UgXl8
|
||||
td+J2j/fUYcZK2qZlces/YOgwQkFx214tLQi7mfb6nS4t1bL53r8tjodHmiJfH3P6uXdMs2zkG2N
|
||||
lGGe5W3NUfunWU1WH3RnM48d9g/sa+qErk0Fvy0uWJveY719OujrmHu5Nm+fDm4JRL5eKjP4ncvo
|
||||
PAvZNhMG51lWNhi1fwbVULDglqbGzDxm7J850AFgglOnTumJJ57QL3/5S6ubApvxraxU+ZxbLFbb
|
||||
uqqYe1rr9iuwqXKeZ65S36ZKDV05oryeA9pw5YHKq//K8INM4/Kkv202GLl/9lVLnScPKK/ngHbf
|
||||
Way2zcYfNBuXx6vtJeE87+i0Z6Mu1lXN8xhkTUWhgt17k7o3OHP3+o4WHX51xTwPLtULr66Q/9hJ
|
||||
7ant0aFjk/Lu2mztQWZGedLf1nAZ7p9ndknv7+jRntoe/eyCR0+/bvFBc0Z5irSuLpznA10tqdbe
|
||||
DpYgQTL9uuIfCn7r3ZvkALBVDXEH7utqAjpY0zjPc7fq+ZrGyDrvHcNDKi7rMfgg08g86W+bOWP3
|
||||
z9Nl0rn+PL1yNk9v+8tVX230QbOReXyq9obz7Nb1gja117Rmo9ELQgeAiXbs2JF0wkHkoirtC515
|
||||
DR+oR39tuBIclq0VlQqXCF/td9R8+4B230n8jBGllWoMTfImSf23b2hIi7Wq0JgkQQbmSXPb7DAy
|
||||
z1U19byl9tAqWtcmHkiexUqzvzlNRubxq703nGdMtyYleZYwCsAUwQN4SdLoqC59Evx2eVuz1v2m
|
||||
Rz+7MM/Dt61UVdSSbp+fGZU/0Zld02SYJ81tjZdpnhEdre1VX+hx90YmpZJCLTWqufPKNM+4+prD
|
||||
eSZ0b1TSykJGASCpu8OHdF2SVK56X0DPxwxmbNXzW7pUH3XgvqxsUNX38vT2fNPZeJu1RkO6FlrS
|
||||
7a7/XY2pXN78rEeIYVieNLfNFuPydOv42QqdCZ2R908NSQWrZPRYVuPy9OtMfzjPNfmnJOWvtmwU
|
||||
AJMAmuzll1/WxMSEfvSjH1ndFFgtPBlb1IF6tP7bodnZVanmUql7ROofeEtNklpXzv3UvsIlku7r
|
||||
4/AyzeP3NSipotAryaB3BgPzSOltmxUG55kROjCfvKFr2Wh3MqblKdYqj6Q79+XuOXQtFL8smyRp
|
||||
Uu9/b2Zpts87e3VU0vqvz/1Uy7/ikTSpe6EDOX0yoTFJxV8pkmTSOu9ZzJPutobIcp4ZpXpml0ca
|
||||
HZWpV38blqdQS0skXZhgSUHMoVvH+1frJV+biiWtqQno4Bxb3x2u0HFJ6+bpJVtWUCFpMHggJklT
|
||||
H+u+Qtd+G7rOuzF50t02e4zLM6NVT5eVS1PvGvUJNooZeULLDfo/TrzUoAkYAZCC0dFRBQKBWV9+
|
||||
v1+jo6P6zW9+E1kOMBWvvfYalwPkvJnJ2GZmY48zfl3vTga/bayx+zXV5FloHl/tJjUqelI+I5iU
|
||||
p3S7Ai3PqfHOO8o7f3WBbUXaRq/oUNRZYscjT0LL22pUJcl//ra1B8zZyLOtTocHNqkqZoUBIImp
|
||||
dr1xNvFZ1uAQ/pkzxY5AnrQsK9urNZLG/CfMOWA2Mo+3Twe3dGmNf7deuWzdwsuMAEjBihWJr2kr
|
||||
Li6O/Ly+vl5/8zd/o//+7//W178+fxf4jh07tH79eq1atSqrbYVT+NXee0DtGW9jF+RZkKJ69VQv
|
||||
liY/VIuhy++ZlGfkhPJGvOpoflGB5iXa0HuOUQBZ9Lv2HlctuUaeNDxZq+d3eaTRKzreac7IDEPz
|
||||
nDqvPaeK1NT7rA73FupQ8wCjADCvS5fzdMnqRmQReVJQ0KFvlZVLU536xbC5nygMyeNv0itnfWrw
|
||||
XdRB32p19LdbMgqAEQBZ5PF4VF9fr4mJCe3bt2/e7V9++WUTWoVc1D9xX9ISrQ5P+le0RBWSBidY
|
||||
490+vOrYvFHlkk5fdtOBcmhZwWQrCsBWPv/9pGKWc3uyUMWSxn5v0vB/pKBITa9Xyyvp6k/ddKA8
|
||||
rkvnJ6WSEq2zemUD5Jy7U4OSKoJDsSWpYLWWSLo/5Z53Y+fzqaE6OBT/+rA1B8rGCE02WLBV1Uav
|
||||
bJAEHQAG8Hg8+tGPfqS33357zu1OnTqlU6dOmdQquFuV+qKXX5t4oCEt1taVwdlLgrO/31CvY0Za
|
||||
xuVxvAR5SjerzSPpzjtqcsx+CYvP41VH8/bI5IFrCxcrZg4K2EipXohefm1wQn55VN0Q7K1Z3lAi
|
||||
r+7okmPemuLyOF6CPNvW6ekSSRc+0FHH7Jew+DxFauqt0/rQT5eWxs1BARimVc9HL782fUtjKtda
|
||||
b/BdbJl3q4p1Wlccc54kLo/jJcjj3af6Akn+3TrumP0SFp/HpwZfn9aFfuotKFfMHBQmowPAQN/8
|
||||
5jfnHQnw7LPPmtQauNtV9d7RzEzy4+fUcuVBZK33i9WLdfqDE7LuaqN0xeVxvPg8wWUaJUkrnlOg
|
||||
Zb8CLfs1WGv0/LbZkmj/VKorlKNrxQN1nnTS/1suGdGlC5qZSf6TAR0/Ninvrmd1eKBFe3d5dPUH
|
||||
5x00ND8uj+PF54madX/zJh0eaNHhgRZ9r80pw2sS7Z8V2hnKsXPzpN7f4aT/NzhXd/DgPjyT/FS7
|
||||
fjE8pOKyizq4JaD2snJdv9zkoCH5cXkcLz5PcJlGSZK3Swe3BHRwS0CQ5T1EAAAgAElEQVQvlTnl
|
||||
tFCi/dOoHaEcO7xDOtdv3f8bcwAY7Lvf/a5ee+21Obf5r//6L33ta18zqUVwg+7zBxIfXN25Ebm/
|
||||
f+At5Q2Y2KgMpJJn3m1tZP48V9XU45xJ8ubP46T5HXJL0uu4L9yO3P95Z6/2dJrYqAykkmfebW1k
|
||||
/jwjOlrrnCFC8+cZV19zj/pMbRVyUdLrt/29kfvvDlfolWETG5WBVPLMu62NzJ+nW8fP2v3T3oz5
|
||||
8/TrTH+ezpjaquQYAWCw4uLieS8F6OnpMak1cK3S7eqSi2ZeJ4+9uS1PLtlWp51y0czr5LE3t+WB
|
||||
c3n7tEPWzryeVeSxN5vnYQSACerr6+f8+ZEjR/T666/L4/GY1CK4zsgJ5bnp8xV57M1teXLJqfPa
|
||||
47jryOdAHntzWx44l79JrzjuOvI5kMfebJ6HEQAmKC4unncugNu3b5vTGAAAAABATqIDwCRf/epX
|
||||
5/z5nTt3TGoJAAAAACAXcQmASVasWDHnz2/dumVSS4Ly8vJM/X2BQMDU3wcAAAAAiEUHgEmKiuZe
|
||||
smdoaMiklgAAAGCh3nrrrYw/t/3hD3/I6PFnzpzR5cuXM3qO//f//p++//3vZ/QcAJyHDgCTzDcC
|
||||
4MiRI/rJT35iUmsAAACwEENDQ7p0ydqF1u7du6d79+5l9BxLly7NUmsAOAlzAAAAAAAAkAPoAAAA
|
||||
AAAAIAfQAQAAAAAAQA5gDoAcxaz8AAAAAJBbGAFgE9u2bbO6CQAAAAAAF2MEgEnu3Lkz5883b95s
|
||||
Uktgml0/tLoF2UUee3NbnhzRcWmX1U3IKvLYm9vywNn+ucHqFmQXeezNbXkywQgAk/zxj3+0ugkA
|
||||
AAAAgBxGB4BJPv300zl/vmrVKpNaAgAAAADIRVwCYJKBgYE5f75kyRKTWgLTHPuh1S3IXPSwcvLY
|
||||
j5vz5Ij2dcesbkLGooeVk8d+3JwHzvaPZ6xuQeaih5WTx37cnCcTjAAwycsvvzznz7/yla+Y1BIA
|
||||
AAAAQC5iBIAJLl++POfPt23bpuLiYpNaE5SXl2fq72PZQQAAAACwFiMATPBv//Zvc/6cJQABAAAA
|
||||
AEajA8Bgd+7c0WuvvTbnNl/96ldNag0AAAAAIFfRAWCwf/qnf5rz59u2bVNNTY1JrQEAAAAA5Crm
|
||||
ADDQL3/5Sx05cmTObb71rW+Z1BoAAABk6vDhwxk/x3vvvaeDBw8u+PEtLS3atWtXxu0AkHsYAWCQ
|
||||
X/7yl9qxY8e82zU2NprQGgAAAABArmMEgAFSPfg/efKkPB6PCS2ajVn5AQAAACC3MAIgiy5fvqy/
|
||||
+7u/S+ng/8UXX2T2fwAAAACAaRgBkIKxsTEVFxfPun9yclJ+v183b95UX1/fvNf7Rztw4EA2mwgA
|
||||
AAAAwJzoAEiB1+vN6vNdunQpYYcCAAAAAABG4RIAk126dIll/wAAAAAApqMDwCTbtm3T6OgoB/8A
|
||||
AAAAAEtwCYAJ3n77bX3zm9+0uhkAAAAAgBxGB4CBuru71dDQoBUrVljdFAAAAABAjqMDIMv27dun
|
||||
rVu3qrKyUh6Px+rmAAAAAAAgiQ6ABXvxxRdVXl6uP/mTP1FZWZkef/xxrVq1yupmweFa6/ary/Oh
|
||||
NvSeU7/VjckC8tib2/K43fqOFu3cnOAHo1d0qHlAn5veosyQx97clgcu4O3TwZrGmLuuX87Tcb9F
|
||||
7ckUeezNbXmiMAmgpEAgkPbXT37yE7300kv627/9W9XX13PwjyyoUvMKSZ5KbS+yui3ZQB57c1ue
|
||||
HFZSrb0DdVpvdTuyhTz25rY8cACfGnyBWQdjkrSmJqCDNa0WtCkT5LE3t+WZjQ4AwCZ8tZsULDWL
|
||||
tXWl1+LWZI489ua2PDnlwgfaU9sT/PrBndCdK7SlzaE9OeSxN7flgcP41OC7qPqC4K2x4Q165Wxe
|
||||
8Ovy6eCd+au1zLoGpok89ua2PInRAQDYglfbSxZHbpWXrJHPwtZkjjz25rY8OezUbV21ug3ZRB57
|
||||
c1se2J93X+Rg7PrlPL0xHHXBmr9JHcNDGvOf0F1rWpc+8tib2/IkQQcAYAdFa7TVI2nyhk5PSvJs
|
||||
1L5SqxuVAfLYm9vy5LJtK1UlSZrUlTPjFjcmC8hjb27LA5vzqaEsNAx7qlPvJ7j2+q7/kH7jd8os
|
||||
NuSxN7flSY5JAAEbaF23UeWShkYv6DUtUWP1YjWurJJGnHmuhTz25rY8OWfzJh0e2DRz2+mTspHH
|
||||
3tyWBw6yVt7wUOxkZ12nunXJzCZlhDz25rY8ydEBAFguNBmbHujd237164aGqjeqfEWlWnVV3VY3
|
||||
L23ksTe35UFwUrZqXf1Bj46esroxWUAee3NbHos888wzeuaZZ6xuhr0VrNaS0Lf3p5x/1pU8Nue2
|
||||
PHPgEgDAaqWVwcnYJm/oxLik8et6d1KSKrW31oGTs5HH3tyWJxdFT8pW+0HkmuyqVx06Mzt57M1t
|
||||
eQAgx9EBAFjKq46aSknS0Oj10Frsfp0YfSDJiZOzkcfe3JYH0oiORmZm92jpk5Y2JgvIY29uywOn
|
||||
WFLgrncn8tib2/LE4xIAwErhydgklVe/qEB13M89ldpedE79obmWfLXf0cXq8OztD9R58i21J5uH
|
||||
qXS7ApsqIzeHrhxRxUCCGU2yycg8YaXbFah5oA2952T4AC0D8ywoe6bMyjP5oTn7B0GDE/Jrhbzy
|
||||
aGmFpE+k5W3N2rsrtLM1qfd39KrvkySP31anw6+uiNz0Hzup1zstnOAt0zxh2+p0+K8nrL9ePcM8
|
||||
C8pupGzmYT4BzGXqhK5Ntam+QCouWCul8K6yrGxQ7WXloVtDOtdfoTNTSTb29sWs9T42vCF21vds
|
||||
MzpPmLdPB8tuqaO/3djZ6g3Os6DsmTAzz1Sn8ftnDowAACzkW1mp8jm3WKy2dVWh76u0r1rqPHlA
|
||||
eT0HtPvOYrVtrk9yxrZKfZsqNXTliPJ6DmjDlQcqr/4rdRi8bLNxeYJa6/bHdGoYzcj9k272bDAu
|
||||
j1fbS8LbvqPTno26WFeVcEsYoKJQwYs3JnVvUJJK9cwu6f0dwWHbP7vg0dOv12p5wgeX6oVXV8h/
|
||||
7KT21Pbo0LFJeXdtVpOVZ3YzyhO0vqMlplPDUhnun3SzGy6jPEVaVxfe9gNdLanW3g6WIEEy/bri
|
||||
Hwp+692rhoJE27SqocwX+f7pMulcf3Dd9rf95aqv7kiyZnurnq9pjKzz3jE8pOKyniS/I1uMzBO0
|
||||
riYQ06lhLGP3T7rZM2dkHp+qveFtd+t6QZvaa1qNCJESRgAAlqnSvtAZ00Rn5yNnVCOTs11VU8/M
|
||||
LO3XJh5IKxYrYR9laaUa9UCdt4PP2X87ONHbqkJJhp3YMzBP6PHNtw9ot/ary5Ngg6wzMk962bPD
|
||||
yDx+tfe+Ffp+TLcmpUbPEvkSbovsCh7AS5JGR3XpE0ka0dHakcgW90Ympc2FWirNPtO6baWqNKn3
|
||||
Q0u6fX5mVP5d1ZEzu+bLMI+CZ5jX/aZHP1OLdq40vsVzyzRPetmNl2mecfU194a+n9C9UalqZaGW
|
||||
J9wWkO4OH9L1si6tUbnqfQF5L+fpeOTtq1XPb+nSGg1J/gqdmerW8bMzU9n6p4Yk7yp5pdlnWr3N
|
||||
WqMhnQst6XbX/67GytrkzZdk4Flmw/IoeIa5+l6e3lZAO/KNyxDNuDzpZbd/nn6d6a8IfX9N/ilp
|
||||
Tf5qLTM4TzJ0AABWCU/GFpqNPV74oL1clWoulbpHon8aOpibvKFrCZ7aV7hE0n19HD7YH7+vQUkV
|
||||
hV5JBl0GYGAeSeofeEtNklpXZrndyRicZ2HbZsC0PMVa5ZF05z4H/0aJX5ZNkjSp97+XaCh1qZ7Z
|
||||
5ZFGR3UvwVMt/4pH0qTuhQ/2P5nQmKTirxTJwN7CWFnMI0mfd/bqqKT1X892Q1OU5TwL2zaLDMtT
|
||||
qKUlki5McPCPOXTreP9qveRrU7GkNTUBHUzpca16uqxcmno34aeeZQUVkgblDx/sT32s+wpd+23o
|
||||
Ou/G5JGku8MVOi5p3dJstTUVxuVZ2LaZMiNPaLlB/8dcAgDklpnJ2CKzsceLzM4uNdbEDr321W5S
|
||||
o6IncrMaeRaax5zsJuUp3a5Ay3NqvPOO8s5fnWtLZNPoFR2qTXzN9fK2GlVJ8p+/7ZyDLPIkZJvs
|
||||
2cizrU6HBzap6sIH2tM+MteWgDTVrjfO5untBEdWwSH8s6+7Xla2V2s0x3ruViJPWnlMz25kHm+f
|
||||
Dm7p0hr/br1y2bqFlxkBAFjCr/beA2pfyDZF9eqpXixNfqgWoyf1Sxl5ItLJY1p2k/KMnFDeiFcd
|
||||
zS8q0LyEiQCz7HftPfpdOg94slbP7/JIo1d03MpJ/ZIgTxp5LMhuaJ5T57XnVJGaep/V4d5CJgJE
|
||||
Si5dztOlVDYs6NC3ysqlqU79wshJ/TJEnhTyWJjdkDz+Jr1y1qcG30Ud9K22bCJARgAAjuJVx+aN
|
||||
Kpd0+nLyg6v+ifuSlmh1eNK/oiWqkDQ4YZcD7LDU8jhHOnmckH0hbQwtK+ip1HaDJ53EXIrU9Hq1
|
||||
vJKu/jT5wdXnv59UzHJuTxaqWNLY7+12gJ1aHudIJ48Tsi+kjeO6dH5SKinROpYTRNb41FAdHL59
|
||||
fTj5wdXdqUFJFcGh2JJUsFpLJN2fstu7cWp5nCOdPE7IvpA2hiYbLNiqakMnnUyODgDASUo3q80j
|
||||
6c47aoq7RruvZb8CzaGh2xMPNKTF2royOFdzcPb3G+q120jLVPM4RTp5km5rIynn8aqjebvC89mu
|
||||
LVysmDkoYL5t6/R0iaQLH+joqegflOqFgRYd7g3N0j44Ib88qm4I9tYsbyiRV3d06dTsp7RUqnmc
|
||||
Ip08Sbe1kZTzFKmpt07rQz9dWho3BwWQKe8+1RdI8u+OmrxNCk7gFtBBX2iW9ulbGlO51nqD72LL
|
||||
vFtVrNO6YrvzJCnmcYp08iTd1kZSzuNTg69P68IPKyhXzBwUJuMSAMAxgkv7SZJWPKdAy3OSwjO6
|
||||
X1XvnefUGJ6lffycWq5U6mLUWu+nP3hL1l1tlEgaeaxrZBrSyTPXtnZ5l0t3/1Sqq2W/uiRJD9R5
|
||||
8oTN/t9ySdRM7VETuPmPndTrnSO6dGGTqsKztH8yoOPHSrR317M6vCv4kKs/6E1v6Lfh0shjXSPT
|
||||
kE6euba1Sw9buvtnhXYOtGinJGlS7+84b7P/NzhXcGk/SZK3Swe3BN+RxoY36I3hbl3xd2lNeJb2
|
||||
qXb9Ynir2ssu6mBZ8CHXL1ekNuTbNGnksa6RaUgnz1zb2uVTYbr7p1E7tgS0Q5I0pHP9TZb9v9EB
|
||||
ADhG7LJsCd25ETno6h94S3kDhjcqA+nlCes+f8CmB5bp5ElhW8ulkyeVOQZgnthl2RK6cDty0PV5
|
||||
Z6/2dBreqAyklycs7WvYTZNOnhS2tVw6ecbV19yjPhNahVwUuyxbQv7eyEHX3eEKvTJseKMykF6e
|
||||
sJSvXTddOnlS2NZy6eTp15n+PJ0xoVWp4BIAwA1Kt6tLLpp5nTz25rY8uWRbnXbKRTOvk8fe3JYH
|
||||
zuXt0w5ZO/N6VpHH3myehxEAgBuMnFCemz5fkcfe3JYnl5w6rz12vY58Ichjb27LA+fyN+kVu1xh
|
||||
lw3ksTeb52EEAAAAAAAAOYAOAAAAAAAAcgAdAAAAAAAA5AA6AAAAAAAAyAF0AAAAAAAAkAPoAAAA
|
||||
AAAAIAfQAQAAAAAAQA6gAwAAAAAAgBxABwAAAAAAADmADgAAAAAAAHIAHQAAAAAAAOSAR6xuAOBa
|
||||
u35odQuyizz25rY8OaLj0i6rm5BV5LE3t+WBs/1zg9UtyC7y2Jvb8mSCEQBAFuXl5VndBAA2R50A
|
||||
kKuof4D16AAAsigQCFjdBAA2R50AkKuof4D1uAQAMMqxH1rdgsxFDysnj/24OU+OaF93zOomZCx6
|
||||
WDl57MfNeeBs/3jG6hZkLnpYOXnsx815MsEIAAAAAAAAcgAdAAAAAAAA5AA6AAAAAAAAyAF0AAAA
|
||||
AAAAkAPoAAAAAAAAIAfQAQAAAAAAQA6gAwAAAAAAgBxABwAAAAAAADmADgBA0qOPPhpz++HDhxa1
|
||||
BIBdUScA5CJqH+AudAAAkgoLC2Nuj4+PW9QSAHZFnQCQi6h9gLvQAQAAcJ38/HyrmwDA5qgTAHIR
|
||||
HQAAANfhgz2A+VAnAOQiOgAAAAAAAMgBdAAAkoqLi61uAoAMTExMGP47qBOAs5lRJ9yI2ge4Cx0A
|
||||
QAJ8SACc5Ysvvoi5bcbQXuoE4CxW1Ak3ovYBzpWfn08HAJBI/IcEAPYW/4HUjA/21AnAWayoE25E
|
||||
7QOcKz8/X49Y3QgAM1rr9qvL86E29J5Tv9WNyQLy2Jub8uTCmb31HS3auTnBD0av6FDzgD43vUWZ
|
||||
IY+9uS2PlBt1wtW8fTpY0xhz1/XLeTrut6g9mSKPvbktTxRGAACSioqKrG6CpCo1r5DkqdR2OzQn
|
||||
Y+SxN3flGRsbi7ltxAd7e9SJBEqqtXegTuutbke2kMfeHJzHjDrhRtbXPp8afIFZB2OStKYmoIM1
|
||||
rRa0KRPksTd35YmvexIdAIAkadGiRTG3E71YjOar3aRgqVmsrSu9pv/+bCOPvbktTzwjPtjboU5I
|
||||
ki58oD21PcGvH9wJ3blCW9qs/pC+QOSxN7fliUIHQGqsrX0+Nfguqr4g9LuHN+iVs3nBr8ung3fm
|
||||
r9YyE1uUGfLYm9vyzFZUVEQHACDN7t1++PChyS3wanvJ4sit8pI18pncguwij725Lc/sD6Rf/vKX
|
||||
s/47rK8TCZy6ratWtyGbyGNvDs9jRp1wI0trn3df5GDs+uU8vTEcdcGav0kdw0Ma85/QXfNalBny
|
||||
2Jvb8mh23Vu0aBEdAIAkPfJI7HQY4+Pj5jagaI22eiRN3tDpSUmejdpXam4Tsoo89ua2PAkYMWTV
|
||||
8jqRyLaVqpIkTerKGRu0J1PksTeX5bF+aLszWFf7fGooCw3DnurU+wmuvb7rP6Tf+J0yiw157M1t
|
||||
eYLiO+yKioqYBBCQZn8I+L//+z9Tf3/ruo0qlzQ0ekGvaYkaqxercWWVNOLMcy3ksTe35ZGk+/fv
|
||||
x9yOH7KaDVbXiYjNm3R4YNPMbQdPyiaJPHbnojxm1Ak3sq72rZU3PBQ72VnXqW5dMqk1mSOPvbkt
|
||||
T1B8hx0jAICQ+A8B5p7ZC03Gpgd697Zf/bdvaEiSVlTKWdOMhJHH3tyWJ2h6ejrmthFn9qytE3Mo
|
||||
qdbegRa9sM3qhmQJeezNwXnMqBNuZFntK1itJaFv708566xrQuSxN7flCUk0AoAOAEAWX99WWhmc
|
||||
jG3yhk6MSxq/rncnJalSe2sdODkbeezNbXlC4q9xKy4uzvrvsM0cANGTstV+ELkmu+pVZ87MTh6b
|
||||
c1EeM+qEG9mm9gFIW3yH3SOPPEIHACBZeWbPq46aSknS0Oj10Frsfp0YfSDJiZOzkcfe3JZnRvwH
|
||||
+8LCwqz/DnuOABjR0cjM7B4tfdLSxmQBeezN2XnMqBNuZIfat6TAqe9OiZHH3tyUJ/6SHeYAAELi
|
||||
e7dNW+ImPBmbpPLqFxWojvu5p1Lbi86pP/Re66v9ji5Wh2Zvn/xQG3rPKekgpdLtCmyqjNwcunJE
|
||||
FQMJZjTJJiPzhJVuV6DmQWrbZsrAPDHb6oE6T76ldqM/U5mVJ9V9mSXT09MxQ3vz8/MNWd7Lsjox
|
||||
n8EJ+bVCXnm0tELSJ9Lytmbt3RXa2fNds72tTodfXRG56T92Uq93Wti5kWmesG11OvzXE9Zfr55h
|
||||
nphtNan3d/Sq7xMT2p1MNvOYOJ+AWXXCjSyrfVMndG2qTfUFUnHBWimFd5VlZYNqLysPPb5THf3t
|
||||
yWds9/bFrPU+Nrwhdtb3bDM6T5i3TwfLbqW2bSYMzhOzrYZ0rr9CZ6ay0fAkzMyT6r7MgvjXa35+
|
||||
PiMAAGn2MECz3tx8KytVPucWi9W2rir0vVfbS6TOkweU1/OOTns26mJdVZLHValvU6WGrhxRXs8B
|
||||
bbjyQOXVf6UOgy93NC5PUGvd/phODaMZuX/2VYe3PaDddxarbXO94WffjcuT/r7MJrOG9VpVJ+ZV
|
||||
UajgxRuTujcoSUVaVye9vyM0ZLukWns7ki3zUKoXXl0h/7GT2lPbo0PHJuXdtVlNVp7ZzShP0PqO
|
||||
lphODUtluH+e2RXetkc/u+DR06/XarlJTU8oozzp78tsYfj/wllX+/p1xT8U/Na7Vw0FibZpVUNZ
|
||||
+N3Tp2qvdK4/T6+c3a3rBW1qr0k2u02rnq9pjKzz3jE8pOKyniS/I1uMzBO0riYQ06lhLGP3z9Nl
|
||||
4W3z9La/XPXVHVqW/RBRjMyT/r7MlkS1jxEAgILD24qKimKGtY2NjRn8AaFK+0JnTBOdnY+cUV1R
|
||||
qVZdVbf8au99K9w63ZqUGj1L5FOCPsrSSjXqgTpvB5+z//YNDVVv1KpCSYad2DMwT+jxzbcPaLf2
|
||||
q8uTYIOsMzLPVTX1zMy4f23igbRisVLrb14oI/Okty+zzawP9tbUifkED+AlSaOjuvSJJI2rr7k3
|
||||
9PMJ3RuVqlYWark0+0zrtpWq0qTeDy3p9vmZUfl3VUfO7JovwzwKnmFe95se/Uwt2rnSjDbPJdM8
|
||||
IzpaOxK5dW9kUtpcqKUJtzVDpnnS25fZRAfAwllZ++4OH9L1si6tUbnqfQF5L+fpeOTtq1XPb+nS
|
||||
Gg1J/gqdmerXmf6K0M+uyT8lrclfrWXS7DOt3mat0ZDOhZZ0u+t/V2NlbfLmSzLwLLNheRQ8w1x9
|
||||
L09vK6AdJg1uMS5Pt46f7Y7c8k8NSd5V8ibc1gl50tuX2UQHADCH4uJic9/cwpOxhWZjjxc+aC9X
|
||||
pZpLpe6R6J8Wa5VH0p37iQ+WC5dIuq+Pw3HG72tQUkWhV5JBlwEYmEeS+gfeUpOk1pXZbXZSBueZ
|
||||
ETown7yha9lpeWKm5Uln2+yIf3MzcmZv0+tEIvHLskmSJvX+9xINpS7U0hJJFyYSHyx/xSNpUvfC
|
||||
B/ufTGhMUvFXimRgb2GsLOaRpM87e3VU0vqvZ7mdqcpynhmlemaXRxod1b3stDQ1huVJZ9vMmVkn
|
||||
3Mi62tet4/2r9ZKvTcWS1tQEdDClx4WWdPN/nPhguaBC0qD84YP9qY91X6Frvw1d592YPJJ0d7hC
|
||||
xyWtW5qlpqbEuDwzWvV0Wbk09a5Rn2CjmJEnnW0zFz9nR2FhIZcAAGHmDnGbmYwtMht7vMjs7FJj
|
||||
TdTw8NLtCrQ8p8Y77yjvvF3WbSfPQvP4ajepUdGT8hnBpDwW7cv4tb2XLFmSZMvM2fIygNErOlSb
|
||||
4LrwbXU6PLBJVRc+0J72kYQPtSXyJLS8rUZVkvznb1s7p0E28liwL82sE25kae2batcbZ/P0doKj
|
||||
v+AQ/rhrw719OrilS2v8u/XK5e7ZD7IaedLKs6xsr9ZIGvOfMOWA2dA8Ju/LZCOfGAEAhJj75uZX
|
||||
e+8BtS9km5ETyhvxqqP5RQWal5g62Vpy5IlIJ09RvXqqF0uTH6rF0AkaTcpj0b787LPPYm4/9thj
|
||||
hv0uKz8E/669R79L5wGnzmvPqSI19T6rw72F1k+GF4c8aeR5slbP7/JIo1d03KQJGg3NY8G+NLNO
|
||||
uJEdOj8vXc7TpVQ29DfplbM+Nfgu6qBvtWmTraWLPCnkKejQt8rKpalO/cLICRoTMCSPyfsyWQcA
|
||||
IwCAkPizAfFnC+wltHSbp1LbE4xi7J+4L2mJVod/VrREFZIGJ4wfPLUwc+dxnlTyeNWxeaPKJZ2+
|
||||
bIdOj7mks3/M35c3b96Muf34448b9rucVSckaVyXzk9KJSVal2Biv89/P6mY5dyeLFSxpLHf22GJ
|
||||
w0TmzuM8qeQpUtPr1fJKuvpTe3V6zJbO/jF3X5pZJ9zIebUvNKFbwVZVJ5jM7e7UoKSK4FBsSSpY
|
||||
rSWS7k/Z9d147jzOk0oenxqqg0Pxrw/bs9NjRjr7x7x9mezSJzoAgJD43u34swXW86qjebvCc4au
|
||||
LVysmev8q9TXsl+B5tDQ7YkHGtJibV0ZnKs5OPv7DfXaatRsGnkcIc08pZvV5pF05x012Wq/hKWT
|
||||
Z65tjTU2NjZraS8zRwDYr05IUpGaeuu0PnRraWn0df6lemGgRYd7QzPJD07IL4+qG4IfCpY3lMir
|
||||
O7p0ypKGJ5FGHkdIM8+2dXq6RNKFD3TUVvslLJ08c21rHLPrhBs5o/b51ODr07rQLW9BuWau82/V
|
||||
81sCOugLzSQ/fUtjKtdab/BdbJl3q4p1WldsdZ4kjTyOkGYe7z7VF0jy746aiM9O0skz17bGuX37
|
||||
dsxtLgEA4pSWxi5F9Omnn1rUkrlUqqtlv7okBdeOP6HgFURX1XvnOTWGZ5IfP6eWK5W6GLXW++kP
|
||||
3pL9rhxLMY+FLUxPqnmCyzRKklY8p0DLc5ISz85vrXT2T7JtjTUyEtt7YvRZPWfUCUlaoZ0DLdop
|
||||
Kbh2/PnQcO4RXbqwSVXhmeQ/GdDxYyXau+tZHd4VfOTVH/SmN/TbFCnmsbCF6Uk1T9Ss+1GT8fmP
|
||||
ndTrJl0KkJp09k+ybY1jdp1wI+fUvkbt2BLQDknBteObQsO4u3XF36U14Znkp9r1i+Gtai+7qINl
|
||||
wUdev1yR2pBvU6WYx8IWpifVPMFlGiVJ3i4d3BL8dDE2vEFvmHwpwNzS2T/JtjVOfO2rqAiuREAH
|
||||
ABBSWFgYs8zN9PS0JiYmVFhYaHHLwlK4jvvOjchBV//AW8obMKFZC5ZenrDu8wds2JEhpZcndhlA
|
||||
e0onTypzDBgj/s1t1apVhv4++9cJKbjUWo/65trkwu3IQdfnnb3a02lCsxYsvTxhaV/Dbpp08sQu
|
||||
A2hP6eRJYVsDmF0n3MgZta9fZ/rzdGauTfy9kYOuu8MVemXYhGYtWHp5wlK+dt106eSJXQbQntLJ
|
||||
k8K2BkjW+cklAECU+B7u+BeObZVuV5fsNIt+hshjbzbKEz+8beXKlYb/TsfWCUnaVqedctgs+nMh
|
||||
j73ZJI8VdcKNHF37vH3aIZvOor8Q5LE3G+SZ69InRgAAUR577IofiisAAAqjSURBVDH97ncz54w+
|
||||
/fRT1dbWWtiiFI2cUJ6D3ofnRR57s1Ge+A+g8R9QjeDYOiGFZl+3uhFZRB57s0keK+qEGzm69vmb
|
||||
9IqdrrDLFHnszQZ55qp7jAAAosSfFYg/awDAPsbGxmJmuH300UdNubaXOgE4h1V1wo2ofYBz0AEA
|
||||
pCj+rMC1a9csagmA+cS/PletWqVFixYZ/nupE4BzWFUn3IjaBzjH9evXY25Hd+DRAQBEKS0tVX5+
|
||||
fuT22NiYTZe6ARD/5rZmzRpTfi91AnAOq+qEG1H7AGd4+PChbt26FXPf2rVrI9/TAQDEWb9+fczt
|
||||
6OvdANhH/Nmn6Dc3o1EnAGewsk64EbUPsL+bN2/qiy++iNwuLi5WcXFx5DYdAECcdevWxdyOP3sA
|
||||
wHqfffbZrOt6n3jiCdN+P3UCsD+r64QbUfsA+5uv45MOACBO/Ivk2rVrevjwoUWtAZDIhx9+GHPb
|
||||
7Ot6qROA/VldJ9yI2gfY36VLl2Jux1/6RAcAEKe4uDhmhuDp6elZHyIAWOvixYsxt//sz/7M1N9P
|
||||
nQDsz+o64UbUPsDeRkZGdPPmzZj7GAEApKCuri7m9tmzZy1qCYB4AwMD+vTTT2Pue+qpp0xvB3UC
|
||||
sC+71Ak3ovYB9nX+/PmY22vXro25/l+iAwBIqL6+Pub2zZs3Z62nCcAaZ86cibn91FNPqbCw0PR2
|
||||
UCcA+7JLnXAjah9gTw8fPpxV+/78z/981nZ0AAAJFBYWzjpT8Ktf/cqi1gAI++ijj/TRRx/F3Jfo
|
||||
zc0M1AnAnuxUJ9yI2gfY07//+7/HzP5fVFQ0a+UOiQ4AIKn4IW4fffSRzp07Z1FrAHz22Wfq6uqK
|
||||
ue+JJ56wdFkv6gRgL3asE25E7QPsZWRkRL/+9a9j7quvr0848WleIBAImNUwwGn+4R/+Ydawtpde
|
||||
einpREJ5eXlmNAtwvfi3pmvXrun73//+rO3efPNNlZaWmtWshKgTgDWcVCfcKN3aJ1H/gGyJrn+f
|
||||
ffaZ/v7v/z7m50VFRfrxj3+s/Pz8WY9lBAAwh+9+97t69NFHY+6Ln0gj2pe+9CWjmwTkpEQf3nfu
|
||||
3GmLD/XUCcAe7Fwn3Cjd2idR/wAjJJrf5IUXXkh48C/RAQDMqbS0VN/97ndn3ZfM/v37jW4S4Hov
|
||||
vvjirPvy8/NjXnv19fV67rnnzGxWUtQJwHxOqxNulG7tk6h/QDbE17/42veNb3xj7pE4XAIAzG9g
|
||||
YEBHjx7VokWL9Oabb1rdHCAnHT16VB999JG+8Y1vzJqF2g6oE4D17F4n3IjaB1jvxz/+sQYGBvTt
|
||||
b3973tpHBwAAAAAAADmASwAAAAAAAMgBdAAAAAAAAJAD6AAAAAAAACAH0AEAAAAAAEAOoAMAAAAA
|
||||
AIAckHkHwFCnNuTlKS+Frw2dQ8HHnN6dfLvdp9P55erckKf0HhJqb1oPspCJf9+hzg0JHzPnnyph
|
||||
+zYo3BQgJdQRAKmiXgDIFHXEWBy/2Jr9RgB0N8X+MyQ1pM4NFWrvT/2pT+/OU15Fu/oltepWzuzk
|
||||
GMn+vkOdaknyx+xuylPehk7F/7mGOjdE/p6x+tVe4aAiBfehjgBIFfUCQKaoI8bi+CW7Apka7Aj4
|
||||
pIDkC3QMpviYvtaApIDUGuiL/UGgVQr9TAFf0ieM3a61L8lmMb8yvH0a7bQDE/++gx2+hPcl3j7q
|
||||
uXwdgcEkj0ll3wDUEQApo14AyBR1xFgcv9iazUYANKorMKgOX/BW/4l3E/fa5DWpO41nHercoKZu
|
||||
SfKpY/Ci2sqz01rnmfvvW952URfj/jjlbRfV15po+0Z1DXbI19qnwMU2lcc8pifyO7p7c6AXDTZD
|
||||
HQGQKuoFgExRR4zF8Uu22awDQJLK1bYvssf0btQeHurcoIr24ACYvkCfWlN5uqihIa192XjxnNbu
|
||||
VK8ZiVzLslunNfsaFWtGmCT/+yZTsTr0ali7KuaFovI2XexqTPg7Vq3NsJlARuxeRwDYB/UCQKbs
|
||||
Xkc4fpl5Ko5fbNgBMLfWvoACgS4l2m2zDamzJXiNh69jUAn3dTpO707Sexe8ZmSu6356d+eFXvwh
|
||||
vg7tzbQ9Jhn8ONhu3+qKFB8xpFvXlOZjAPNYWkcAOAr1AkCmOH4xH8cvydm8A2CtVkV12ZS3XUzv
|
||||
RTD0rk70S1Krtqsls5kehzq1oSn00vF1aDAQUCD0FRli0l6RpFesW93dwRdx+DHxw06sEfv3TSQy
|
||||
/Ki1b9bwmqROHwpNbuLT9q3Wp0Sus1EdAWBz1AsAmbJRHeH4heOXBLLYARCaOTHjJRWG1Pla+B91
|
||||
tTLpfxl690Rohsdutc+aIXL+Xq9opw+FZov0dWgw7p+/sWvmRdT92uzZJoMPG0z9HzAhc/++p3cH
|
||||
n7+iXeoYDCiQcuWKev7WfQx9RJrcXUcAZBP1AkCm3F1HOH7h+CURm40AOK3deTNLY7Tuy6SXaUjv
|
||||
noh60bT2zfReRU8k0d6Swj/gafVG/icSt6mxOdyN9rEGZ/3ULj1JC/n7hl+4weuA5jPU2RJ6/lb1
|
||||
MfYRlrBrHQFgP9QLAJmyax3h+IXjl8Sy2AHgC/a0BOK/5pq4oltNMb1BM9enZPXaOV+HBmOerFxt
|
||||
F8OTcPTrxHwzSQzdUvCSEJ+SXhJSsVrB1+Q13Zr1dPMPVZmfuX/fxq7Q84e7BtWtpvlm/YiasMTX
|
||||
sTfF65yAaC6uIwCyjHoBIFMuriMcv4jjl8RsNgJAketTMhtuEveU27cm6ClqVKTTK8FyHa6V7t+3
|
||||
sWvmRdT92hy9jTMTlqR1vQ1gBOoIgFRRLwBkijpiLI5fssriDoBW9cX3CJk4uURkeQjXytLft7E5
|
||||
0tv48ezxQZJyb+gM7IQ6AiBV1AsAmaKOGIvjF6PZbwRA1sy/lmN4eYj5n2qVgk+V/B9Igx+HJuzI
|
||||
xnAZh4lZqzTVJU4AJ8hiHQHgctQLAJni+MU0OXz84uIOgJmJLRIPkZmZGCPxEJuYZ4oMt+nuTXQd
|
||||
SfTMkc3u+wc63Ru69ibRNUSntbtiZuhMDnWeIUdkr44AcDvqBYBMcfySJRy/JOXqDoDI0I/+dlXE
|
||||
TAAxpM4NTZF/ilRmuIzMktndpLxZz1URWTeyY69T/4OG1LkhwZIcMeuHblf8n+r07vDfMbeGziCH
|
||||
ZLGOAHA56gWATHH8kgaOXxbC3R0AalTXYEdwdsvupqjZJKOWkuibaxbK6Kfq0mB47Y0kz+Xr6HHu
|
||||
upGnD6m9P8FaneHeMbWqL/76m9O7FX5tzZ6xM+prQ+K1RQFnyGIdAeBy1AsAmeL4JWUcvyyIyzsA
|
||||
JJW36WLUupkzghNMpNPpU952UYFAePmN2c/l5Jkjhyr2zhSIOL6OQQUC8dfGnNbumVcP4G5ZrCMA
|
||||
XI56ASBTHL+khOOXhckLBAIBqxsBAAAAAACM5f4RAAAAAAAAgA4AAAAAAAByAR0AAAAAAADkADoA
|
||||
AAAAAADIAXQAAAAAAACQA+gAAAAAAAAgB9ABAAAAAABADqADAAAAAACAHEAHAAAAAAAAOYAOAAAA
|
||||
AAAAcgAdAAAAAAAA5AA6AAAAAAAAyAF0AAAAAAAAkAPoAAAAAAAAIAfQAQAAAAAAQA6gAwAAAAAA
|
||||
gBxABwAAAAAAADmADgAAAAAAAHIAHQAAAAAAAOSA/x88lETUMs7ifwAAAABJRU5ErkJggg==
|
||||
"
|
||||
id="image63"
|
||||
x="0"
|
||||
y="-1282.8716" />
|
||||
</g>
|
||||
</svg>
|
||||
|
After Width: | Height: | Size: 27 KiB |
Reference in New Issue
Block a user