<table border="1" cellspacing="0" cellpadding="8">
<tr>
<th>Issue</th>
<td>
<a href=https://github.com/llvm/llvm-project/issues/104520>104520</a>
</td>
</tr>
<tr>
<th>Summary</th>
<td>
[AMDGPU] Masked load/store with a uniform but non-trivial mask don't simplify
</td>
</tr>
<tr>
<th>Labels</th>
<td>
new issue
</td>
</tr>
<tr>
<th>Assignees</th>
<td>
</td>
</tr>
<tr>
<th>Reporter</th>
<td>
krzysz00
</td>
</tr>
</table>
<pre>
[godbolt link](https://godbolt.org/#z:OYLghAFBqd5QCxAYwPYBMCmBRdBLAF1QCcAaPECAMzwBtMA7AQwFtMQByARg9KtQYEAysib0QXACx8BBAKoBnTAAUAHpwAMvAFYTStJg1C1aANxakl9ZATwDKjdAGFUtAK4sGexwBk8DTAA5DwAjTGIQAGYATlIAB1QFQjsGF3dPPQSk2wE/AOCWMIiYy0xrHIYhAiZiAjSPLy5S8pSqmoI8oNDwqNiFatr6jKb%2B9s6Cot6ASktUN2Jkdg4AUkiAIQBqAFkMN3oASQARDdXjsA4AOgB6FiYFAGtMdABaVDiFC5Nz5Y0AQQU5gtMAB9Gj0ZhsE6RY7LABMsOutweT1e70%2BtDhsJ%2BvwGwEwBA26CY1QMAE85gTTid4ZhnnEQAA2SSMyR0rgsll02FRbmRWF0yI8nl05lMzlxACsQr5dIZ0v5cQA7BIGRoQLCJXKNXKZXEABwSWEGrhGunRCTRbna9Wa4V4MVM56mLhyl1O2HMmWmPnC0ySA2O0zRV2m52Wm0M9227VOiUmkBx/nOjQeiQp1mmWEaf3q7N654MH2OoQy34S54AcS4BYoyoN0Ux2Nx%2BI2BGIeDi9ChMPhrHQwGQDGefeHLHQCAUTEbf2xq02ADE3AwbCkNji2woQBsWG5%2BnFiKhgMRMAoFBsGKgqMfMOeSJhkPMlLeFKTl7elwB3fzoDZfkzHgh5gYbdMBYEhSQgGpgDYFgt2PJh0A/dsCEwKZsSwGgAg2fdUBQmwng2PtgDiNxgUeYgAloDZTFQPAfwAASRR50GBfw238JJkAUCA4jbQj0HQYgFDiJhFggLgplvURePmG8LyXDCNng9ABFoUlCNoPBgGAl1qQlVRSGwviEME4TRMwcTJIvaTAOPd8GEUpDCEwVT1LELSdIZPTSUk2hUFEWhgSXCEnmBEziGpSINA2MAwGPABHFiPxIe5gSPOY4mBJIAC8b1i6LlkVNZsUYNtSRAbENj0uju1bJg6A2AKqPkhzMCoJTDDxCA8D5DYNEMk09UknrYQ2eiTHMC4iMHC5kuIe5nJYC46IuVQIDQv4qrhCUVty1RKWhDY9oJBgAmADYRuqn8iAuplKr0vyEIuNAHNqvBkBYOINj2AlLu2uiBoZe7tse9BnoEMGkiMegiAYC5THvN6GCUWoylA0qoScSQNlUC6uFWbBsNogEGEMvAuAe1Anpe9AyaZPrgY1UHwYc%2BHEapBQEDcKgqHoBGbBITHsdx8mCcp6mIYuKHgBhgQ2eQQzVixnG8bFhI8BJxXImVkW%2BTF3KD38ZJ3INxmJWZ3iIqpPEUPoNhBEtoWVZGgnDMd8KzLEiS9IM27sf%2BsHjrN0wxDcG8qWqBqmqd3WsUiQnxrMJamKeT4qbBv0RouOIuB44yBKEkSvckkH0%2Bztsyd6l0tZ11X4/FsGaflmvhYuvX64N2iGGNzSDY234to1fo7xZn8qXez7vtoX7eoDwy9c2vTh%2BPUepf8GX8Tl/mkZR230cEGO68J9XNbxpeiBXmm6exjQzeXzBV%2B3jmuZ5vn7wvw/Rfr7b79X6XZbhvzFuzt8b1xPgIYBsd9bhC7j3PAptF4/wvg/R21t8Ro3tgQVB2tW4u3jm7fOpki4WW9ttUk18roXCDovSOVFo40RqonSaKdIbIPhpILOOcIBK1we3Qm20Q7uEwAQiKHtiGWXPiPS2ldRrV0/qA/hQ82FN35v3KqAFqK0XQLORUhxZwznWBsRcy4KhrgIBuLcO49wHiPCeM81kxC0BCKJe4t4ryYDkokV8yB7JflesJe8exiRMBCF2P8tAAJARAmBYgEELwBH7lgZABg7ItUUpdJhS1ppwzmgtFCS0VprRLrCSI%2BMZwGIXEuFcAgzEWO3LuLBNjjyniko45xyBXEXncZ4l8b4Wp%2BJ/OEyJFFongUgsQaCoE4KYAQok%2B8KTw44OdnwsaE1k53GYmnJ6mc%2BTZ1zu7AunsSFWX8kwGSx4ZEXRYLcCZkCj53LwdgYpC9fizkMcY6pwF1xCUsQ0nCtiWkOJMO0zpl5rzPm8b478v46ARPxFEmCYyoIwS3E5FCczkk1BvAwhiayLgsKlmwnZCIuE8OWXHbAIj%2BJEPMhIhx5zhFt1Gnga5UEHkKOeZEfRvxiTsRCG4FCZ44RRVqoVTYVjGmHmafYsFHjbzHgfEJHpkL%2BnQqGfCkZiLYnjMmbBJSMzELIVQtSWERESLPF5sSZ4uTnjpTcHEZ4OVMCNmhJiQapBBqYhNWatwBZUDDmIsQL1mIfV%2BueGgT6MMUjDk%2BcG3sY5iK%2BovM8DCTAfrPHimHMOcbTUJvNcm/AZkCDIAQM8OiObQ0Fo1iJYtpbLYVrzUm/1CAZkOvrfCb1jaw0IESMWxxdI2wNv7Pm/1LLOzvUIFBAdQaO0hq7cm2g6AFDPHIpRMt2jZ3xuHU254O5p5aXbC8XpyBhwTKHYmsNmbMBh2neekd1qUp2odXRZ4qg707tyU%2B9dzxSTvrDZ%2Bg89rv3ZT/cm3Ji1v1vs3bm7d/6UoQZfb%2B6Dlb/XgfycBnNS48D8GICwB981bWAYdU6l1ej4RtmzbO3R2JeXtn5YKyKFMqRivqdYqVdjWnApcW48FF5j1Qv8XEQJBhqihJvOq2ywEtVxIEMawqei/i0bwPRk8kVRrMaKqxyVAKZVNRBTxuVfGVVzAYAMmF/4NVSdAkiiZKL9WzJONRxT5i6MCtU8KyIorNMSv%2BdKzjTjuNdN414vpJmzMSYRdZ7VyKpm/iNSXJzry/j5VqrFMVl13VMo2E0LL%2BNEscBmLQTgEpeBeG4LwVAnAfA%2BAAGpbA2PsAAShsAE8xFiMd4AQTQBWZj3B5BcWE0QNB6j1JECUiomTZk1PoTgkhSvddIJVjgvBNz9S6xwLQMw4CwCQBGuIdBwjkEoHtg7EQTA2GIEue4fA6AoR%2BRAEIC2Qj%2BBqOVcrpBnvMFiQAeRCNod%2BnAeCkAjZg77DA1ILawPy4AThHGbne1gW40MlhaFIPgBVtgEbw9R5gVQgSUKA94GxMoC3NIhGIK9lwWAFvsVgu9hGxAQiJEwIcUCnVNJGG6zMS1wAFA1bwJgD832hNlaB/wQQIgxDsBNDIQQigVDqA2zoPQBhOdrP0MpzckAZhvAqPD3d7FOzOuhNk0cLw%2BwTiYBsXdyASKnGAFQVQ0RJCwgqwzw9J54AzCsO/FIDgHJDEaKQXw/guiFB6E0LIyQBCB8yIkaPDBxjdAiCMNGny2iDFcA0PQPv08DA6KHiYEfLD59jyMfPSfw8p%2B94CRYEhCvFfm0rxbnANjnY2BACjDB7iSQgLgQggthUSU61z3rURJAXEiIqcbo2NBDdhAyWEJoJQzY4HN0gZXUdLZWyANbXPSDbcQCACkJECBHYgCd%2BgxBAisCWO3zvl3u9TF4E8Afh69Di%2BEAFaXsJZfyCUGoAtroE0EhGcoTg3hwCVhvgtktt9gKqfhsJeG3rQD4g/ldr3i4J9Kdoxs/qQOtptmPpEJEJPkQaQWQWQavuvpvhVpwDvnvkrlMBAa7tAc3tvngaPqQAzkkPYJIEAA%3D%3D)
The suboptimal codegeneration I'm seeing relates to `llvm.masked.load` and `llvm.masked.store` intrinsics where the mask is uniform across the vector but not a constant true or false. That is, for example, in cases where each lane will either do or not do the masked operation in its entirity. This _should_ compile to a vector operation, but does not.
That is, given the input IR
```llvm
; ModuleID = './masked-ops.ll'
source_filename = "./masked-ops.ll"
target datalayout = "e-p:64:64-p1:64:64-p2:32:32-p3:32:32-p4:64:64-p5:32:32-p6:32:32-p7:160:256:256:32-p8:128:128-p9:192:256:256:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-v2048:2048-n32:64-S32-A5-G1-ni:7:8:9"
target triple = "amdgcn-amd-amdhsa"
; Function Attrs: mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite)
define protected amdgpu_kernel void @masked_intrinsics(ptr addrspace(1) nocapture noundef readonly align 16 %x, ptr addrspace(1) nocapture noundef writeonly align 16 %y) local_unnamed_addr #0 !reqd_work_group_size !0 {
entry:
%id = tail call noundef range(i32 0, 128) i32 @llvm.amdgcn.workitem.id.x()
%id.zext = zext nneg i32 %id to i64
%load.cond = icmp ult i32 %id, 16
%load.cond.singleton.vec = insertelement <4 x i1> poison, i1 %load.cond, i64 0
%load.cond.vec = shufflevector <4 x i1> %load.cond.singleton.vec, <4 x i1> poison, <4 x i32> zeroinitializer
%load.ptr = getelementptr <4 x i32>, ptr addrspace(1) %x, i64 %id.zext
%value = tail call <4 x i32> @llvm.masked.load.v4i32.p1(ptr addrspace(1) %load.ptr, i32 16, <4 x i1> %load.cond.vec, <4 x i32> zeroinitializer)
%store.cond = icmp ult i32 %id, 32
%store.cond.singleton.vec = insertelement <4 x i1> poison, i1 %store.cond, i64 0
%store.cond.vec = shufflevector <4 x i1> %store.cond.singleton.vec, <4 x i1> poison, <4 x i32> zeroinitializer
%store.ptr = getelementptr <4 x i32>, ptr addrspace(1) %y, i64 %id.zext
tail call void @llvm.masked.store.v4i32.p1(<4 x i32> %value, ptr addrspace(1) %store.ptr, i32 16, <4 x i1> %store.cond.vec)
ret void
}
; Function Attrs: mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none)
declare noundef i32 @llvm.amdgcn.workitem.id.x() #1
; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: read)
declare <4 x i32> @llvm.masked.load.v4i32.p1(ptr addrspace(1) nocapture, i32 immarg, <4 x i1>, <4 x i32>) #2
; Function Attrs: mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: write)
declare void @llvm.masked.store.v4i32.p1(<4 x i32>, ptr addrspace(1) nocapture, i32 immarg, <4 x i1>) #3
attributes #0 = { mustprogress nofree norecurse nosync nounwind willreturn memory(argmem: readwrite) "amdgpu-flat-work-group-size"="128,128" "amdgpu-no-agpr" "amdgpu-no-completion-action" "amdgpu-no-default-queue" "amdgpu-no-dispatch-id" "amdgpu-no-dispatch-ptr" "amdgpu-no-heap-ptr" "amdgpu-no-hostcall-ptr" "amdgpu-no-implicitarg-ptr" "amdgpu-no-lds-kernel-id" "amdgpu-no-multigrid-sync-arg" "amdgpu-no-queue-ptr" "amdgpu-no-workgroup-id-x" "amdgpu-no-workgroup-id-y" "amdgpu-no-workgroup-id-z" "amdgpu-no-workitem-id-x" "amdgpu-no-workitem-id-y" "amdgpu-no-workitem-id-z" "uniform-work-group-size"="true" }
attributes #1 = { mustprogress nocallback nofree nosync nounwind speculatable willreturn memory(none) }
attributes #2 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: read) }
attributes #3 = { mustprogress nocallback nofree nosync nounwind willreturn memory(argmem: write) }
!0 = !{i32 128, i32 1, i32 1}
```
I get assembly (at `-O3) like
```assembly
.text
.amdgcn_target "amdgcn-amd-amdhsa--gfx90a"
.amdhsa_code_object_version 5
.protected masked_intrinsics ; -- Begin function masked_intrinsics
.globl masked_intrinsics
.p2align 8
.type masked_intrinsics,@function
masked_intrinsics: ; @masked_intrinsics
; %bb.0: ; %entry
v_cmp_gt_u32_e32 vcc, 16, v0
s_load_dwordx4 s[0:3], s[4:5], 0x0
v_cndmask_b32_e64 v1, 0, 1, vcc
v_lshlrev_b16_e32 v2, 1, v1
v_or_b32_e32 v1, v1, v2
v_lshlrev_b16_e32 v2, 2, v1
v_and_b32_e32 v1, 3, v1
s_mov_b32 s4, 0
v_or_b32_e32 v8, v1, v2
v_lshlrev_b32_e32 v1, 4, v0
s_waitcnt lgkmcnt(0)
v_mov_b32_e32 v2, s1
v_add_co_u32_e64 v6, s[0:1], s0, v1
s_mov_b32 s5, s4
v_addc_co_u32_e64 v7, s[0:1], 0, v2, s[0:1]
s_mov_b32 s6, s4
s_mov_b32 s7, s4
v_pk_mov_b32 v[2:3], s[4:5], s[4:5] op_sel:[0,1]
v_pk_mov_b32 v[4:5], s[6:7], s[6:7] op_sel:[0,1]
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_2
; %bb.1: ; %cond.load
global_load_dword v2, v[6:7], off
v_mov_b32_e32 v3, 0
v_mov_b32_e32 v4, v3
v_mov_b32_e32 v5, v3
.LBB0_2: ; %else
s_or_b64 exec, exec, s[0:1]
v_lshrrev_b16_e32 v9, 1, v8
v_and_b32_e32 v9, 1, v9
v_cmp_eq_u32_e32 vcc, 1, v9
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_4
; %bb.3: ; %cond.load1
global_load_dword v3, v[6:7], off offset:4
.LBB0_4: ; %else2
s_or_b64 exec, exec, s[0:1]
v_lshrrev_b16_e32 v9, 2, v8
v_and_b32_e32 v9, 1, v9
v_cmp_eq_u32_e32 vcc, 1, v9
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_6
; %bb.5: ; %cond.load4
global_load_dword v4, v[6:7], off offset:8
.LBB0_6: ; %else5
s_or_b64 exec, exec, s[0:1]
v_lshrrev_b16_e32 v8, 3, v8
v_and_b32_e32 v8, 1, v8
v_cmp_eq_u32_e32 vcc, 1, v8
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_8
; %bb.7: ; %cond.load7
global_load_dword v5, v[6:7], off offset:12
.LBB0_8: ; %else8
s_or_b64 exec, exec, s[0:1]
v_cmp_gt_u32_e32 vcc, 32, v0
v_cndmask_b32_e64 v0, 0, 1, vcc
v_lshlrev_b16_e32 v6, 1, v0
v_or_b32_e32 v0, v0, v6
v_lshlrev_b16_e32 v6, 2, v0
v_and_b32_e32 v0, 3, v0
v_or_b32_e32 v6, v0, v6
v_mov_b32_e32 v7, s3
v_add_co_u32_e64 v0, s[0:1], s2, v1
v_addc_co_u32_e64 v1, s[0:1], 0, v7, s[0:1]
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_10
; %bb.9: ; %cond.store
s_waitcnt vmcnt(0)
global_store_dword v[0:1], v2, off
.LBB0_10: ; %else11
s_or_b64 exec, exec, s[0:1]
s_waitcnt vmcnt(0)
v_lshrrev_b16_e32 v2, 1, v6
v_and_b32_e32 v2, 1, v2
v_cmp_eq_u32_e32 vcc, 1, v2
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_12
; %bb.11: ; %cond.store12
global_store_dword v[0:1], v3, off offset:4
.LBB0_12: ; %else13
s_or_b64 exec, exec, s[0:1]
v_lshrrev_b16_e32 v2, 2, v6
v_and_b32_e32 v2, 1, v2
v_cmp_eq_u32_e32 vcc, 1, v2
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_14
; %bb.13: ; %cond.store14
global_store_dword v[0:1], v4, off offset:8
.LBB0_14: ; %else15
s_or_b64 exec, exec, s[0:1]
v_lshrrev_b16_e32 v2, 3, v6
v_and_b32_e32 v2, 1, v2
v_cmp_eq_u32_e32 vcc, 1, v2
s_and_saveexec_b64 s[0:1], vcc
s_cbranch_execz .LBB0_16
; %bb.15: ; %cond.store16
global_store_dword v[0:1], v5, off offset:12
.LBB0_16: ; %else17
s_endpgm
.section .rodata,"a",@progbits
.p2align 6, 0x0
.amdhsa_kernel masked_intrinsics
.amdhsa_group_segment_fixed_size 0
.amdhsa_private_segment_fixed_size 0
.amdhsa_kernarg_size 16
.amdhsa_user_sgpr_count 6
.amdhsa_user_sgpr_private_segment_buffer 1
.amdhsa_user_sgpr_dispatch_ptr 0
.amdhsa_user_sgpr_queue_ptr 0
.amdhsa_user_sgpr_kernarg_segment_ptr 1
.amdhsa_user_sgpr_dispatch_id 0
.amdhsa_user_sgpr_flat_scratch_init 0
.amdhsa_user_sgpr_kernarg_preload_length 0
.amdhsa_user_sgpr_kernarg_preload_offset 0
.amdhsa_user_sgpr_private_segment_size 0
.amdhsa_uses_dynamic_stack 0
.amdhsa_system_sgpr_private_segment_wavefront_offset 0
.amdhsa_system_sgpr_workgroup_id_x 1
.amdhsa_system_sgpr_workgroup_id_y 0
.amdhsa_system_sgpr_workgroup_id_z 0
.amdhsa_system_sgpr_workgroup_info 0
.amdhsa_system_vgpr_workitem_id 0
.amdhsa_next_free_vgpr 10
.amdhsa_next_free_sgpr 8
.amdhsa_accum_offset 12
.amdhsa_reserve_vcc 1
.amdhsa_reserve_flat_scratch 0
.amdhsa_reserve_xnack_mask 1
.amdhsa_float_round_mode_32 0
.amdhsa_float_round_mode_16_64 0
.amdhsa_float_denorm_mode_32 3
.amdhsa_float_denorm_mode_16_64 3
.amdhsa_dx10_clamp 1
.amdhsa_ieee_mode 1
.amdhsa_fp16_overflow 0
.amdhsa_tg_split 0
.amdhsa_exception_fp_ieee_invalid_op 0
.amdhsa_exception_fp_denorm_src 0
.amdhsa_exception_fp_ieee_div_zero 0
.amdhsa_exception_fp_ieee_overflow 0
.amdhsa_exception_fp_ieee_underflow 0
.amdhsa_exception_fp_ieee_inexact 0
.amdhsa_exception_int_div_zero 0
.end_amdhsa_kernel
.text
.Lfunc_end0:
.size masked_intrinsics, .Lfunc_end0-masked_intrinsics
; -- End function
.section .AMDGPU.csdata,"",@progbits
; Kernel info:
; codeLenInByte = 404
; NumSgprs: 12
; NumVgprs: 10
; NumAgprs: 0
; TotalNumVgprs: 10
; ScratchSize: 0
; MemoryBound: 0
; FloatMode: 240
; IeeeMode: 1
; LDSByteSize: 0 bytes/workgroup (compile time only)
; SGPRBlocks: 1
; VGPRBlocks: 1
; NumSGPRsForWavesPerEU: 12
; NumVGPRsForWavesPerEU: 10
; AccumOffset: 12
; Occupancy: 8
; WaveLimiterHint : 0
; COMPUTE_PGM_RSRC2:SCRATCH_EN: 0
; COMPUTE_PGM_RSRC2:USER_SGPR: 6
; COMPUTE_PGM_RSRC2:TRAP_HANDLER: 0
; COMPUTE_PGM_RSRC2:TGID_X_EN: 1
; COMPUTE_PGM_RSRC2:TGID_Y_EN: 0
; COMPUTE_PGM_RSRC2:TGID_Z_EN: 0
; COMPUTE_PGM_RSRC2:TIDIG_COMP_CNT: 0
; COMPUTE_PGM_RSRC3_GFX90A:ACCUM_OFFSET: 2
; COMPUTE_PGM_RSRC3_GFX90A:TG_SPLIT: 0
.text
.p2alignl 6, 3212836864
.fill 256, 4, 3212836864
.section ".note.GNU-stack","",@progbits
.amdgpu_metadata
---
amdhsa.kernels:
- .agpr_count: 0
.args:
- .address_space: global
.name: x
.offset: 0
.size: 8
.value_kind: global_buffer
- .address_space: global
.name: !str y
.offset: 8
.size: 8
.value_kind: global_buffer
.group_segment_fixed_size: 0
.kernarg_segment_align: 8
.kernarg_segment_size: 16
.max_flat_workgroup_size: 128
.name: masked_intrinsics
.private_segment_fixed_size: 0
.reqd_workgroup_size:
- 128
- 1
- 1
.sgpr_count: 12
.sgpr_spill_count: 0
.symbol: masked_intrinsics.kd
.uniform_work_group_size: 1
.uses_dynamic_stack: false
.vgpr_count: 10
.vgpr_spill_count: 0
.wavefront_size: 64
amdhsa.target: amdgcn-amd-amdhsa--gfx90a
amdhsa.version:
- 1
- 2
...
.end_amdgpu_metadata
```
(`-global-isel` acts similarly)
## Expected code
I expected that, with optimizations on, there should be one conditional guarding a `global_load_dwordx4` and one guarding a `global_store_dwordx4` - instead, the masked vector load isn't vectorized.
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJzMXF2TqrjT_zTuDaUFARQv9iIgKCoqvuuNhRAhypsQQP30TwHqOI7OzNndp-o_dY4j9C_dnXSnuxPCGHGMbR-hvyu8WOFbfxkJcYLo70N0OccXmv5rG1jnnGYH1jZwCeVi_1DhWxUgOISEcYWFFaBUgHKl14LILm6wlwoLh6u-7UBFPFq8Lp3gKhutRE3yxLEl9qHC6JIJjZEsQe2SiUSDDahnCtF08QzHdrNH9JUMzzHe0voSSidBFGEvEP0phDPYCTOopVBZTSekazMSY8DByTi4zTWcZq3e3oJtZUZgj4vb6JSJB6E1hZBvZXA_bas6bK_g1FWhyOiKvo_bCmuNRvrkADK5AhQ4lBaaildn-hR11JUD8RrD_WTUP_MTIZwcvUAVgkErOw6wYpCovu9tK4AHYjOuSwGpw8mBzAage7BsDs4OGKq6eITKQRuwAZzocNyylvsYcnBoi3UFrxTYJpoFRWPeworEDy78ioa6PuNtokGx2d7TayeW6-NVow9FLR4mJENTBjXoQihpOXG3-JZm7UzORAjq0orRNSgLvD2dYotbO74v6xCCyUQ7j-nIjl2XBsp4640XCn1WBd-leVe7uCcoxXrEW-qaTnlThw1RbY8DvS91l73BsrdeylDMJgu5HTntxB9L0_HWN5oLg5tz2lzj617fObvMEHQ0b-GNfIk-TyAY0pfxHHg8WHi01gSgMR_imXcQTxWg-MO9jrAeMp63kI0de2wM6jyndcAw0M8sN-E5aE44cRWcgwE9O4HBqeipCsS2ilq4KTnayIk6etbWW220Q51-R5fgbCpvd-B0pAFsySzMttJhsMcgC3QxXrT5goWv4J5ja2MNBooYt3v2UdtpQzTpOoeR5vaR0pu6DeRmIru7BK317nDp2A5vr7aWJq5kZ6Lb7dBurRSbgFHXofNJYMWSnPXW8SRr29AGO2sm6l7m22BE7BYenOwZWtnQDVrrqaKP4C6DcDIe83Rb3GWAFbrdA5xJkNtv9T3d0Ve20sJdR7Htvh266WyMRl5b6J-XLWng200oKsHMlQbBSLTb0GWOujsfT9pZimVN5qZjLTOnXTU1pnA4soS2PAtbkpzNp0xf7k8sdT2aTmYH4MwUeeHYk6Uk-54oX3A7nKgDCDSYtUdQ7Ch4dFIRb09GwxXniV0VSsKWwfW-e5gP1vFsNSCTMdzKgz2d7BNm27XrEMDecXRwLpoUdKet1kmCQotvrQaydqCbs4MfrXSA8LSjSRzWOhJ_SFTE-92VxNXVZK6rLT3ljpEjzebkzIx7C6e1agZd0YYebK3GydEXMExCtxfR6VlWkwHsDI6peBBXQ3WwB9DtJyDBYrBGDRKjpugHstY-YA3ZGK4kftpJB_W2NFsEZxfSx2BiTvSBO5PqibLYivEgsOW2rKUwgYiBfoia8GysR5G9YmadbGUWvtORjVDU5ZbZs4-dQGxvRXXaia2T4ElmWPfUVdLr2KJj6wDp-mWOT35b2CqOKojd01L15odFwWaqsL0WK6y7rDo4ed3FmsXqMRzJs1EwcESZBIv5um37Lc0Sepe-jCdNVQONeFc0jjv7aEBnp5bZFo6Lo3j0jha7iGdY90-R1jW8Hppyx62Y0uMgGaqJyHFnsScfJql5OHTK0HEh8dmquzRZM8ySqwBFidvGzvVSZ5Wkyzo3ALjdHlwmrRVgORIwu6BxWvvCcYkujZQEywYKWtm5aYJ06hQMRzMU7oT2Upi6FaB4h_UYsLQpt4d1odtcDranEGFx6Xl1dNrrF7S8iPOiHbv3kjWf-o0szSpAGe-ixrbRXNaX662TXpTkQoQtcxrZ6kqMrebWkRp7dgRDouRKZ2mecMaAIcI4YImtz0VAFlzC7r3GbhgeMLdYNAmZHRghCpZSK0gnw7kScPT4GPgTo2epW3XkhGp_ODRVscdkiNU9QOuNRM6gjnsdgtuL5QiPJsC1xtGcrgDlmM93RxeAMmD5lO0doXLs0Uu9PxzPnNM6u_gLMR6bZ67nRLYqJn1zyDWzDh5JSBOYw0miRanXRRzhovOqdQgOc-EsNro7F8VdJKATHmuipICRQCDswjH0xJUty_1M7DRcew1Fu6GS2Sy0us6ECcPs4tlEkSb0PIzG_UiaqMWoauusrS7lRDGhfZEXQ7af9MXBfn_28SHg-NNZXMpL31xzrrDlFmHRplsBylA-dxUl8O2ZHeuGFMhcbwX14FAAhN4047ih5WexITMHnm2vPL9b98qZMV0zl6QBEz-eKIulcMjakxO3lpfjVd-U1gdpz7P9zLJXodTc4ctBM1f1MFOYkzSLddqXCF4chl1tSF9Cl48uI0_Ycg0h5epwPDopsietzoGaTViZVyTzcjgwYgozBxvnvh33XC_z15KayMJwsexsoboveySPNEd1TvzFUQjT9m2Dn8mdQ2-I1vJunJ7wNB5LbU6R1hw3Hs-PwXy6mu_bXsczdqtY6ey3l469k7ud8NAvM4d-bO-kwxr3V_5e87Zid6Jd5L4618lkIY_lSX-k8PFkzSVTfx9OzY685oZzTRp4ThfrRiiKx5UzWBgkE9ezVstTY73PeOqunAoJatvTuqbN9q1ywM02G6943CsuglQTWtIU7qJRB3sdbWhup7NuOjC9vGxZH6SWCnVNGOLhVqeVgMHdFQm3_dU8usxyCCPB9WGYRjtpy-yNjFblibY4nSwVZIO0MSog_f7wks5URYSWbgxp59gwPWDXodIaddRwrBGwX7OnpDMDPDfkQ95YbiVhGZ5FJ5v2OisvHnhbTXSAj9DBYo6zEATWcrzmjnaDbpDzrCyVrMtpMk2j1rq4tIVmlLLraQUopsd2MZMyacxuvYZVAUp9FujrXVli1VvutgKUpb07m6I_mHBCS2irUiY2BWa7gCtrVncZtOfGxLtsh-wYdNGS-Pyqd-4p2nGM06O1Os6t_cKHobngBEXh94yeCktZ74rOETvddHgE50wm85Mqa0Z2RLgCFExvs9H4IBzm40gzesfTeQ576_lgLE5PyXzXns8vcAU1Ly-EucF8YsEDxutev7m4dIeD8WnK47omkZnQO6tBNNAmadEbq3fZT_fJwBpwDJeGXe-iTVZja92Yn3thbgw8mCy7l-hcAcqenx8jRZu7DDvXAiDT_dIdB3Jsil5fn9pyiETU5vrZHPpwBwsqhO1wC8EohnDixFAWTgrWLA7qU11crPesyIvmQqb1gVjfM11ZmuHAm64PzCiY2gdnklqW2p3v7VQcXJqT-inr61rGSZmki-1EtcRsf5CDDrEbsj6FoK3KCeJ2sA6DcSF9DGVVAno5EcXjcG5GoeinmoJ1iMbSILDKGQWRdID7td1otCCj9o1sdObg1Om0tz5i7AVHa65jKOrIFufElrdZc8zbc92ezHqwybVPWiCxA9Vpq72mmy2gksZzO2k67RPUcVfOVHMm-fNBt92uX7QJk0EFMtusa7cEFjgD1xDzJKLrqn1qxeKgpeo6tuetIwSXYKSLztCKRrQWXswDXDspPI5arNVQhpdh4gwOtNnPij5kU7wS46V1hpOe34aKPtfpcVfK7GN2yZpdnoMXKR7mOVZt2Z1uSzZ6-jJPdnaCFbEu0331rIuSKJzVgzFqiae9BbGkDdr-GeCWYkJ6pI9SWmjBeq-DV_Ku3-Kb-7O2G012mVAG7yYnjVeyI6XKzqOj09aHg72_AiKU9i2xgSe63bd12ybObDKrMxc2mo27cClj9WhuAzAQW5YDzwu5G3T1y4rrDGh7vZxZ_fXiDAfb5WjS2pq6BKVDV5bEzsqWSuvZJxHiWFoM2UvqsklzqnCyAHdOvdnCBUCGirGMu-vdWZq1A0iiQKZlp20GU5t1MmnupDY5ENK0e8edE3dRm430FpcnddeKWMx1mz0rOMV52tWHJCTeKJS7cnd00A19oS90I03SMHXmYdZK_fQQuZoIjQaBJktS3zZGRx1eDofRqqvKMJ8VbOv60azQrQoNy8-pg6g42QYhwZ7hUmZgIRv5KDIIDnxKrYCGR8UIYd-mIuQaBMUUCahKnXbd1Kt5RnxAVs0NDKtSpynDt55JMQkilNOwTyLsx9iMqcxBEaKIg6gcROGYSny8CyKPMswoiOOClCKTBBG1TQjlB4QyKDPwY2L4hCJRgqggonaGG6MaNXUMQuG4AiRqF0QUOhle6KL8EvuUacToJhAZpkO5ho-oDLsuhTBxUERZQc4rF2EFd52QRQXhbRCwT2ESU8gnOMLknEvEMbWJnSBxrQ1lBl6IXZQPi3HT-t44VyPvghWgOBdS-zz2d81tnCK_EI_9MCGUOr4C63T5Lx_T6y1WpLTASlyktqgK26IqoFHL685C72oQxjXXrYBGiY6DJDLRZodd5BseujYALxqAsgExIhsRyjKI4RrnICG3JqgaVlhY54qPasg8XoAKC9nyoxqyjxfcI4x_pNQfLxoVFjJ1usJCwNfvnzlFyCng-lkNm_mXJvgKxDcpKVMvmFVTwJWk9CYm5YQrplm_MkzvzKrpnWM15Zn8Ls-AasrQBZ_8VzUFdMEi_1X1C7Z1rjphQRXy1TZT9XGFhXlfclDzeUxJhEP3bgLDs2zTrxqelf93YuMO_7Czkvhm4YOQkCiusJDykpiEUWBHKM4dahchRPlBhMwkivNv8dk3KT9I_Az7VuHoESJJ5FMe8oI8owpGZHvIy3lFyLCyCBN0jwkW2mEfUWEUEGQSZFG5kmGyOaDIRy6VBtiiKhxdes7mY0pXgBCSiDIsK4pDw0QVIDAV0KT8wDRCkkSoUMlCu0Jm4LtnynCx7VNMnaoA_pRPgV8yKBT-wuGcg93ANNxN4ueObm1yXlQFsDRVAUyEjtYmC6LDxo6CJNzE-ILy-zRVaYhl35FPonOFvY5-zjTvLNuiiIFdyjRc96MThm_nKmIWUHSue-6goEnl1xWujH-leWu5TEyQV8NW7VQpYFcBpYTaBZ3KKVZ88X1kl2wK8SSgcr--a5QH2poZ-KVi2PRCKnHJR4NCl_qDgHuDWox920Uk8GspMsvmfowiglzkIT9XQeKoE4WZCitTYYDjMnhh5hOb4lado-iXMm6cYyfZ7Vx0DYafOL_XKWf9Ronb7XzCydQFRQH2McGGiy8oetYk96NcCRvd-lbeeWDx1t1uvph38cM8DxJSw03Qk1d81u5m_4fUWEs5zIJayLybJg-qF9JZkFvxeTyeh_rnkfnka0Um_sl5WPCyxb_zng8-r9znQcrv_OedWv_Sge6c_6UDnd860IfL3ALpl0rp0VWe_OrqfN-Jvuv_vRd9HvEHL4kQKVS7JqBG64_zUd67rWEePlLT54QUh8hMXIMYWxe9zE5-4D-mI9M1HmL_7-JrHvSZ_1zzn1PpF7X_fWC4p7-bQbHnGZH9bNQvLn4dhT-vJ_7dKDwXE-Uw_Lmz_1wO_DQeRffZx-4bhER4m-QLmLIoyKuwhvj_VlDdSrwwqe5cg1RzR60WxUc1Lz7yYi9fjIGidJCKT_DQxg-qhh1GX27mqw0X5TasGma5wnhCWGhnJC6pHhOUR4tnKo5Dg5hONY_372hFBHkiOsgIXxOCmORu85KIvdDFJs4L4Jd014qrZXH5SiEvcQm2I2xVcytUC0s_QYpevmSdD3g53tiqnr4nn78nX16S87DznveN-pr1jXrjfF0Av_WSfM1bQG8x-bM7M-_c-b-Kx-8Eg38q-Ffh9J1U9v9H6sfUfcp8xTqhWLUxlYZYJNZi1pY59uPLvdlt1V5eqnkpQRlxjLyte6ZymYSq1OnqkC2WLfiAntrdsDctmjVyLyPyqzL5ba7ryleLyWrV3p2a9OOismjmxMbGDCy0CbZ7ZJJNiqI4Twj8B-q--qvQzS8rPar8yZNJtUqJyMY-tbtlla_rwjtT2w227iuGD3JBsaKr0E3hod_nEL1sBqQKR98kl_ivIBYWmr5cst6TYgXw222NvoMBXy4FbzqkG9MLNzbZJCzYIBZQqWleF1pAolL6Dow3eVLfWFkQWSeOiiu8mHNli3MjUnHNVVjIX6_pE_0ow7dyHTfbXEido9LCr8rVZSHHNB_Qbuy4EUo3W6ZeqgQ-gMwDLohKhjnkRs0_wY-8wBdehm89MWOfMPHGC9IcQ8Vcqf1rTYSfNPkkhfsyzJmBiekTyrUPnumTChDoh23NZnpT46E78aeeWNbGDEp75kNdv5knNxdzMxf9vnd8AeA-szQ_8Wy84klfu_xMeiGi_iTigdT4Ij083KlphRfBN173eE0F4SZGboWFhTJA-qTMF7bPbOrFZtfX6-_ZxoUjxUaK0AmZm22d-zJMj74eb8xtZPims8nhF6rWF0V6A56nL5NP3x9-ruhi4VNsWd9k5MHJcB9m79VG6VMng93unZOxz_7-iVq6MPuOzD-Sbx18CEdujB7GI59JdY7KhyNvd_v90qGKORV9mt3Nj0ghvJvdD5jmUxxExy9x8An3nxiYezYw-w8MzHxnYfaNhfP_MSIVFnKPBuGeDAL-Y4uA_3WL1J8twv8Di3DfWYT70SLCo0XqTxbh_zuLCB_Z7a1FhNfz6DuLCP-tRYRnizT-gUUa31mE_9EiDHg0ifBkEuEfm-R1ycWCp1rgRd1E_1HdVP8AvqtW6Cs1_6z_yOurhp-8hv7wrHfy6m_kfUoZZQ3AflPU0C-Lmhdl3XPlwryvXL4WNf-pQzP0s0c3Hxyq3PEtnid_LQXTV4Xg1aOLNjeXflYKfErsdz1-nEgPXs4w_9DNf1D_RWB6KPPr71zsAQN-GZjAf2zHr-XZL-qzL2ZmwB9Zkv0-gRdPdP_AqOx_l04eFlT_y1b7UnMxvyi6vlrtS47_1mrc90me4f7Mav9hEQA-QvX_stW-1GXMLwqzr1ar_5HV-B8KgeIMxh-YrfHQS-Rboe19bAHF6LrH06xFgWUQowKkCgDF1laxCxRGgb3F5OVWUpFMH3dabptg1yMN7_ejHrDXMwPI9pBPNjt8QlZ5goB-hQ4jnBoE_Rqfa2JEdol4sMMDJIlRtIntMNqYQeIT6gfQswbbZLdDEcV83-q2_b8JSfRa0w9sse3-G-C9c1dV8ia_1QNbP3HfuQbZxGZUwn1MfqtOGKGiznWRbxPnT1uVTv9Tq2crvPeAJEbxxjr7hofNTUwM8_AaF59jgrzX3DMjRbso8Mm32j1yuD_i2GBrc3ptlbf48x_yv_wR3t8F3-LTGx7nV-_8xEcnstlFCBV4ivkBlCtBCa8whmkm3m1cH6qSB0iEYhSlaJOa5uuRvAEeXfa12jfkyTfMw6Y4IvmS484NDLKJgsS3Nl5goU1xGuk3QKa--TiE8QprIT-IvDtX9nfIku1LsHVi6I3pGl74ui8YIVQwedPVkKlvghRFOzfIXitO7E0cuu8iADqZKMyzyGYXlsKwnxoutjZB-IsW117Gkflb9hZONxcUvXHjr_jvO_cVn_jWnzXAPjoZ5o_Dg33ySvca8q3Np8z5Qfp4NFXr7xLfzBM4fT9IV2RwfHn3EId6aFN9k4x_qCM-FRTVKiX7FvX5wdBTFQG1Vns0q5nxRzHxrpZgRapXFgp5TProFCsWR6X7yFd98UzKI2Ec_VBBDxJvYoflIYvH9dAg8eb3-_Sn-_B2_-H2NCCG-67NpAwjk3x4PzfTiqebYj7lnyhKPm21wCpaAO6BoiKEbgTm43a_Ncl7eBdCbc8ExRWg3OM1VQHC_RQ09hAV-O75YyGbK9oejUU3MA_xE_P5O0I-fO3ROFaCaGGkKB6hSJ69GsvXoIduwTx4D29l6qf2Q9NMQsM3zznhYTctZ9bHHiYo6uDidNunIZSG2mg2lTejtrYZT8ZSvrCcSGM4lTobefAL8Gwijzd593Js_XvsdAxHmw4ctPry-Besp221tVle1WB-gV39TuUCu_4lVm2p7U1O2UiD6Q8N2E1bWTZzGoSSNNM2Q0WZyEUr8KtW0_ZmMuqrn-R8eWB-XRC4VL3cRmSAwNaF-sdStbbDrksBvn5_4vgK9BFEKgDU_ICgWnswqxYl2zWEvI8l1-f2YbLxEDGK2FMQqtVq-aWMr7UyvsYfR5GrVM241_63dRX9ERtrRmQ_wIufvI1lRSiON-UJqrxZubK7wmq-4aHPq7TT53hbuy_u7gj6CRFfo4LwdL84qrg54DL85D_XVWW5FPlJ00-K3nh-1bcCmJhE1PlHtW_q3fWlvpD-TPPauyXhgxcW7J7XP-W69GHEviBuXD7OcVNUzTNOZeX4USTfceChC_dB-iaV1t4vUB-1r93PzH-S99l0n4QXN15f1uIvDsyAB5UKchxi1_0APQ5jfPa2gfuyY7WD9QC8HqR6Puv_EA2p2tflVk4uXid64JR-1fg-NOkP6n4sxm7Cb1HkOsfLkzs54f3JnUf89ZjOw_g_DHX1FilrtU_vFz3UbV-jztMZpQoQKnW6Wvp6FcfILV7oMklMxdjDrhE9pvRrE7YCWEo-heW7InlBVFJUCt1uEscgeTTNMHGo4i0zfCnejIqp8mg2Kd7NKl-morZ57YAoM_AtnIMMl7ITI7Kwb1MGVanTX55Rnbjbi2d5w5fgh42sEl2lsB-T4oiZ9PjK1_W8ec6cwrFfAQ1yvYcvyLoO7V_W36zVZJvGX-hvpgHYRp1jAP-X87fRFGjAcaC5a263NG9aDZ7fGrTAcVvE083GX_hvQAOOFhgesADQfA00hIbRoIFQZwXDqqMKRyPPwG6tOCkbRPZfOI4T9DdDczyg_3KNLXLj4g_WAOCjjCqoebbhW39Ff-eNqtvEjisc7eKYxB9sCCZu8Zduysq3wrcorexycQwCKMUQlTYy7i_jla_e-VUS4RQb5VYZZQXluMTFyc7d-a8kcv9--pM4mDjJtmYGXgUoxXtr5a9qGAV7ZJIKUArN8xry2rX0b_B_AQAA__-M5ZE_">