From b30c7e00d41e29151040aa177d888a1a677e67ae Mon Sep 17 00:00:00 2001 From: George Hotz <72895+geohot@users.noreply.github.com> Date: Wed, 29 Jul 2026 12:01:59 -0700 Subject: [PATCH] support 2d on UNSHARD (kimi) (#17285) * support 2d on UNSHARD * fixes * Fix test and spec * single barrier * 2d sharding works for devices too * cleanups * no _rewrap --- extra/gemm_fragment.py | 116 +++++++++++---------- spec/tinyspec.pdf | Bin 98084 -> 98812 bytes spec/tinyspec.tex | 12 +-- test/backend/test_multitensor.py | 43 ++++++++ test/unit/test_multitensor.py | 2 +- tinygrad/callify.py | 2 +- tinygrad/schedule/multi.py | 172 +++++++++++++++++++++---------- tinygrad/uop/ops.py | 30 ++++-- tinygrad/uop/spec.py | 6 +- 9 files changed, 253 insertions(+), 130 deletions(-) diff --git a/extra/gemm_fragment.py b/extra/gemm_fragment.py index c1d5900692..efa229c3e6 100644 --- a/extra/gemm_fragment.py +++ b/extra/gemm_fragment.py @@ -2,10 +2,12 @@ tilelang-style matmul_relu written with tinygrad UOp APIs. Demonstrates that tilelang's T.alloc_fragment is expressible with existing -tinygrad primitives: a per-thread REG buffer, wrapped in Ops.UNSHARD over a -LOCAL thread range to form the full logical tile. The kernel is written -against the full-tile UNSHARD view, and multi_pm (the same pass that lowers -multi-device UNSHARDs) resolves it into per-thread shard code. +tinygrad primitives: a per-thread REG buffer, wrapped in one Ops.UNSHARD per +sharded axis over the LOCAL thread-grid ranges to form the full logical tile. +Here the 64 threads are an 8x8 grid and each thread owns an 8x8 sub-tile -- +the 2-D fragment layout tilelang infers. The kernel is written against the +full-tile UNSHARD view, and multi_pm (the same pass that lowers multi-device +UNSHARDs) resolves it into per-thread shard code. Reference tilelang kernel: @@ -30,14 +32,15 @@ Reference tilelang kernel: API mapping (tilelang -> tinygrad UOps, idioms from test/backend/test_custom_kernel.py): - T.Kernel(gx, gy, threads=T) -> AxisType.GLOBAL ranges (blocks) + AxisType.LOCAL range (threads) + T.Kernel(gx, gy, threads=T) -> AxisType.GLOBAL ranges (blocks) + AxisType.LOCAL ranges (thread grid) T.alloc_shared(shape, dtype) -> UOp.placeholder(shape, dtype, slot, AddrSpace.LOCAL) - T.alloc_fragment(shape, dt) -> per-thread REG placeholder, wrapped in Ops.UNSHARD over the AxisType.LOCAL - range: fragment.unshard(axis, tid). The full logical tile is shard x threads, - exactly like device sharding, but the sharding axis is the thread axis - carried by the RANGE metadata instead of a device tuple. C_local[i, j] with - i in this thread's shard is INDEX on the UNSHARD, which multi_pm resolves - into INDEX on the per-thread REG shard. + T.alloc_fragment(shape, dt) -> per-thread REG placeholder, wrapped in one Ops.UNSHARD per sharded axis over + the AxisType.LOCAL ranges: fragment.unshard((axis_y, axis_x), (ty, tx)). + The full logical tile is the shard with each sharded axis multiplied by its + range size, exactly like device sharding, but the sharding axes are thread + axes carried by the RANGE metadata instead of a device tuple. C_local[i, j] + with [i, j] in this thread's shard is INDEX on the UNSHARD, which multi_pm + resolves into INDEX on the per-thread REG shard, axis by axis. T.copy(gmem_slice, smem) -> smem[thread_idx].set(gmem_slice[thread_idx], end=copy_rng). set returns the smem tile AFTER the copy; the implicit-barrier pass turns the store->load dependency of the loop that consumes it into a workgroup barrier @@ -51,7 +54,7 @@ API mapping (tilelang -> tinygrad UOps, idioms from test/backend/test_custom_ker from tinygrad.dtype import dtypes, AddrSpace, DType from tinygrad.uop.ops import UOp, Ops, AxisType, KernelInfo -from tinygrad.helpers import cdiv +from tinygrad.helpers import cdiv, getenv from tinygrad.tensor import Tensor # --------------------------------------------------------------------------- @@ -62,35 +65,37 @@ def alloc_shared(shape:tuple[int, ...], dtype:DType) -> UOp: """T.alloc_shared: one LOCAL buffer shared by all threads in the block.""" return UOp.placeholder(tuple(shape), dtype, next(UOp.unique_num), AddrSpace.LOCAL) -def alloc_fragment(shape:tuple[int, ...], dtype:DType, axis:int, tnum:UOp) -> UOp: - """T.alloc_fragment: per-thread REG fragment + UNSHARD over the LOCAL axis. +def alloc_fragment(shape:tuple[int, ...], dtype:DType, axes:tuple[int, ...], rngs:tuple[UOp, ...]) -> UOp: + """T.alloc_fragment: per-thread REG fragment + UNSHARD over the LOCAL thread grid. - Each thread privately owns shape[axis]//threads elements along `axis` in a - REG buffer. The UNSHARD over the LOCAL thread range presents the full - logical tile: full_shape = shard_shape with `axis` multiplied by the number - of threads. This is exactly how UNSHARD carries the DEVICE axis today, - except the sharding axis is the thread axis carried by the RANGE metadata. + Each thread privately owns shape[axis]//threads elements along every sharded + axis in a REG buffer. The UNSHARDs over the LOCAL thread ranges present the + full logical tile: full_shape = shard_shape with each sharded axis multiplied + by its range size. This is exactly how UNSHARD carries a DEVICE axis today, + except the sharding axes are thread axes carried by the RANGE metadata. """ - nthreads = int(tnum.vmax) + 1 - assert tnum.op is Ops.RANGE and tnum.arg[-1] is AxisType.LOCAL, "fragment must shard over a LOCAL range" - assert shape[axis] % nthreads == 0 - shard_shape = tuple(s // nthreads if i == axis else s for i, s in enumerate(shape)) + assert len(axes) == len(rngs) + assert all(tnum.op is Ops.RANGE and tnum.arg[-1] is AxisType.LOCAL for tnum in rngs), "fragments shard over LOCAL ranges" + assert all(shape[a] % (int(rng.vmax)+1) == 0 for a, rng in zip(axes, rngs)) + by_axis = dict(zip(axes, rngs)) + shard_shape = tuple(s // (int(by_axis[i].vmax)+1) if i in by_axis else s for i, s in enumerate(shape)) fragment = UOp.placeholder(shard_shape, dtype, next(UOp.unique_num), AddrSpace.REG) - return fragment.unshard(axis, tnum) + return fragment.unshard(axes, rngs) # --------------------------------------------------------------------------- # GEMM kernel: C = relu(A @ B), float inputs (fp16 or fp32), fp32 fragment accumulator, no WMMA # --------------------------------------------------------------------------- -# 64x64 output tile per block, 64 threads (tilelang uses 128 with a 2D per-thread fragment layout; -# an UNSHARD splits one axis, so here each thread owns BLOCK_M//THREADS full tile rows of BLOCK_N fp32) +# 64x64 output tile per block, 64 threads as an 8x8 grid; each thread owns an 8x8 fragment sub-tile +# (the 2-D per-thread layout tilelang infers for this GEMM) BLOCK_M = BLOCK_N = BLOCK_K = 64 -THREADS = 64 -ROWS = BLOCK_M // THREADS # fragment rows per thread -ROWS_B = BLOCK_N // THREADS # B tile columns copied per thread +TY = TX = 8 +THREADS = TY * TX +TM = BLOCK_M // TY # fragment rows per thread +TN = BLOCK_N // TX # fragment columns per thread def matmul_relu_kernel(c:UOp, a:UOp, b:UOp) -> UOp: - """C[M, N] = relu(A[M, K] @ B[K, N]) -- one 64x64 tile per block, locals + fragments.""" + """C[M, N] = relu(A[M, K] @ B[K, N]) -- one 64x64 tile per block, locals + a 2-D fragment.""" M, K = a.shape K2, N = b.shape assert K == K2 and a.dtype == b.dtype == c.dtype and not dtypes.is_int(a.dtype) @@ -99,53 +104,56 @@ def matmul_relu_kernel(c:UOp, a:UOp, b:UOp) -> UOp: # with T.Kernel(T.ceildiv(N, BLOCK_N), T.ceildiv(M, BLOCK_M), threads=64) as (bx, by): bx = UOp.range(cdiv(N, BLOCK_N), 0, AxisType.GLOBAL) by = UOp.range(cdiv(M, BLOCK_M), 1, AxisType.GLOBAL) - tid = UOp.range(THREADS, 2, AxisType.LOCAL) + ty = UOp.range(TY, 2, AxisType.LOCAL) + tx = UOp.range(TX, 3, AxisType.LOCAL) # A_shared = T.alloc_shared((BLOCK_M, BLOCK_K), dtype) # B_shared = T.alloc_shared((BLOCK_K, BLOCK_N), dtype) A_shared = alloc_shared((BLOCK_M, BLOCK_K), a.dtype) B_shared = alloc_shared((BLOCK_K, BLOCK_N), b.dtype) - # C_local = T.alloc_fragment((BLOCK_M, BLOCK_N), accum_dtype) - C_local = alloc_fragment((BLOCK_M, BLOCK_N), dtypes.float32, axis=0, tnum=tid) + # C_local = T.alloc_fragment((BLOCK_M, BLOCK_N), accum_dtype) -- an 8x8 REG tile per thread of the 8x8 grid + C_local = alloc_fragment((BLOCK_M, BLOCK_N), dtypes.float32, (0, 1), (ty, tx)) - # T.clear(C_local) -- each thread zeroes its own fragment rows - ir0, j0 = UOp.range(ROWS, 4, AxisType.LOOP), UOp.range(BLOCK_N, 5, AxisType.LOOP) - C_loc = C_local[tid*ROWS + ir0, j0].set(0.0, end=(ir0, j0)) + # T.clear(C_local) -- each thread zeroes its own fragment sub-tile + ic, jc = UOp.range(TM, 4, AxisType.LOOP), UOp.range(TN, 5, AxisType.LOOP) + C_loc = C_local[ty*TM + ic, tx*TN + jc].set(0.0, end=(ic, jc)) # for ko in T.Pipelined(T.ceildiv(K, BLOCK_K), num_stages=3): # (num_stages pipelining is async copy + multi-buffering; this is the synchronous single-buffer version) - ko = UOp.range(cdiv(K, BLOCK_K), 3, AxisType.LOOP) + ko = UOp.range(cdiv(K, BLOCK_K), 6, AxisType.LOOP) - # T.copy(A[by * BLOCK_M, ko * BLOCK_K], A_shared) -- each thread copies ROWS row(s) + # T.copy(A[by * BLOCK_M, ko * BLOCK_K], A_shared) -- each thread copies its own 8x8 sub-tile # set returns the tile AFTER the copy; codegen turns that store->load dependency into a workgroup barrier - iar, ka = UOp.range(ROWS, 6, AxisType.LOOP), UOp.range(BLOCK_K, 7, AxisType.LOOP) - A_shared = A_shared[tid*ROWS + iar, ka].set(a[by*BLOCK_M + tid*ROWS + iar, ko*BLOCK_K + ka], end=(iar, ka)) + iar, ka = UOp.range(TM, 7, AxisType.LOOP), UOp.range(TN, 8, AxisType.LOOP) + A_store = A_shared[ty*TM + iar, tx*TN + ka].store(a[by*BLOCK_M + ty*TM + iar, ko*BLOCK_K + tx*TN + ka]).end(iar, ka) - # T.copy(B[ko * BLOCK_K, bx * BLOCK_N], B_shared) -- each thread copies ROWS_B column(s) - kb, ibr = UOp.range(BLOCK_K, 8, AxisType.LOOP), UOp.range(ROWS_B, 9, AxisType.LOOP) - B_shared = B_shared[kb, tid*ROWS_B + ibr].set(b[ko*BLOCK_K + kb, bx*BLOCK_N + tid*ROWS_B + ibr], end=(kb, ibr)) + # T.copy(B[ko * BLOCK_K, bx * BLOCK_N], B_shared) + kb, ibr = UOp.range(TM, 9, AxisType.LOOP), UOp.range(TN, 10, AxisType.LOOP) + B_store = B_shared[ty*TM + kb, tx*TN + ibr].store(b[ko*BLOCK_K + ty*TM + kb, bx*BLOCK_N + tx*TN + ibr]).end(kb, ibr) - # T.gemm(A_shared, B_shared, C_local), no WMMA -- per-thread accumulate over its fragment rows. + # get the shared after the stores (single barrier) + A_shared = A_shared.after(A_store, B_store) + B_shared = B_shared.after(A_store, B_store) + + # T.gemm(A_shared, B_shared, C_local), no WMMA -- per-thread accumulate over its fragment sub-tile. # identical to custom_gemm: a self-referential store over the loop-carried kk range, # which codegen turns into a register accumulator - # kk nests outside the fragment row/col loops (lower ids nest outer), so B loads are contiguous and - # the A row value is loaded once per kk - ir, kk = UOp.range(ROWS, 10, AxisType.LOOP), UOp.range(BLOCK_K, 11, AxisType.LOOP) - jj = UOp.range(BLOCK_N, 12, AxisType.LOOP) - acc = C_loc.after(kk)[tid*ROWS + ir, jj] + A_shared[tid*ROWS + ir, kk].cast(dtypes.float32) * B_shared[kk, jj].cast(dtypes.float32) + ir, kk = UOp.range(TM, 11, AxisType.LOOP), UOp.range(BLOCK_K, 12, AxisType.LOOP) + jj = UOp.range(TN, 13, AxisType.LOOP) + acc = C_loc.after(kk)[ty*TM + ir, tx*TN + jj] + A_shared[ty*TM + ir, kk].cast(dtypes.float32) * B_shared[kk, tx*TN + jj].cast(dtypes.float32) # closing the ko loop here too; codegen adds the barrier so no thread overwrites the tiles while others still read them - C_loc = C_loc[tid*ROWS + ir, jj].set(acc, end=(kk, ir, jj, ko)) + C_loc = C_loc[ty*TM + ir, tx*TN + jj].set(acc, end=(kk, ir, jj, ko)) # for i, j in T.Parallel(BLOCK_M, BLOCK_N): C_local[i, j] = T.max(C_local[i, j], 0) # T.copy(C_local, C[by * BLOCK_M, bx * BLOCK_N]) -- per-thread store of the fragment shard (relu fused into it) # LOOP: these loops are the per-thread output layout; convert_loop_to_global must not globalize them - ie, je = UOp.range(ROWS, 13, AxisType.LOOP), UOp.range(BLOCK_N, 14, AxisType.LOOP) - c_st = c[by*BLOCK_M + tid*ROWS + ie, bx*BLOCK_N + je].store(C_loc[tid*ROWS + ie, je].relu().cast(c.dtype)) + ie, je = UOp.range(TM, 14, AxisType.LOOP), UOp.range(TN, 15, AxisType.LOOP) + c_st = c[by*BLOCK_M + ty*TM + ie, bx*BLOCK_N + tx*TN + je].store(C_loc[ty*TM + ie, tx*TN + je].relu().cast(c.dtype)) # all open ranges are closed at the final store (ko was closed above). # the fragment UNSHARDs go to codegen as is: multi_pm there resolves the full-tile view into per-thread shard code - return c_st.end(je, ie, tid, by, bx).sink(arg=KernelInfo(name="matmul_relu", opts_to_apply=())) + return c_st.end(je, ie, tx, ty, bx, by).sink(arg=KernelInfo(name="matmul_relu", opts_to_apply=())) # --------------------------------------------------------------------------- # python wrapper: same signature as the tilelang function @@ -163,7 +171,7 @@ def matmul_relu(a:Tensor, b:Tensor) -> Tensor: if __name__ == "__main__": from tinygrad import Device assert Device[Device.DEFAULT].renderer.has_local, "this GPU-style kernel needs a backend with local memory (LOCAL ranges + barriers)" - M = K = N = 256 # 4x4 grid of 64x64 tiles, 4 K chunks + M = K = N = getenv("N", 256) # 4x4 grid of 64x64 tiles, 4 K chunks a = Tensor.randn(M, K, dtype=dtypes.float16).contiguous() b = Tensor.randn(K, N, dtype=dtypes.float16).contiguous() diff --git a/spec/tinyspec.pdf b/spec/tinyspec.pdf index 8f358d4e823bc9fec62fec498e4f10148d1f59c6..4e380f5d9db2f68a7b1fa922f8b1c353bd40f2f9 100644 GIT binary patch delta 19543 zcmZ5{LvY{^&}@>8H@0otwr$(m*uU7eZQIyvY}>YzjkE9nRlURa-eFF2n5sE+^>p`? zd_#N&Lnd$n9P9!Dur997=EnB0o*UPC((!oh=zXV}Pro3AH9Ht*JP|-qesqq{)6W>| z)M~(uHqzsc?asb*Bo9~J%uBU_4N2It*IWLx^W|9tGU6wR~&Yz`_<3^Us=??1oYBF13PW}v+0 z5QdK2t}sz3Vp+!AbH+se6mTSnkGhA#@#2%AIO9&tkA6%YR7~tIEkD!MA@&bu)>%Hq zUKrtL8;w!bUgp>NKssHVX!$wK6k~w|@n27e2mPz}O0j&QN9i%;B{HAspPX76_S&e? zkiam|Ypq`MH+~g>pgj(J9JZZ~CDm|_?>08*9#rla_HvjCNwcyKDfqqYEFKjwwC9Qj(R>IA_p-yr{jckOdN3Vi?BTp2- z+bu<(!XpHSanq8P2~Jmx`W0&U4FZX6jK&X42?-kJ9PtSbCbv*QDL5T~BTKfE>!IX& zRz0|xS-;0BkTFWHO-NaE=Qib9RE7JeNTsFsNSx$Oh41z7L#$dbE0L?r>nbAbC+P7g z6AVBQrn&|~g#K-k^V#tJTi|q@LRQpbtkNb{Y31Q&^9!Mhz@zu$<3s4W#~kb1-AN@7C$_6)%MM%_)IGe^LDmHnQ!_!`}6g| zbopMD83&^11P()^_a9=AX3FGYQM-o=<0M4oCl0>0Yvs1mJ}JK7N0?)vGgFAy1`3L# zTf`fPSfhFeLTzd_)0Y!{)$$lM0#$$tw-y$6vYiMzPl*`_Mn6=$hUtEDR#diwoH$g? zGgtFRa6T>kku4c1@4C(5-;j~)t3UR|)QD2mbwxJR-@NJ!&60&J;nLQd!e<2KxN9*P z*dAj+X2zHDxSNG3U?lJDLT3ZZR7IWIrq>0uRWl6nj5%`3FAyYd<;o;w(fI&7Qh(z@ zqO#Kw3dG5H|F&L7RgK686yNjBF;%S2vB&biHbjg=HKaY`CN&Ed3_G;2(ng)Us0c+;LsFPrgFybJ1YTX7RNqd}8wlIsYW%EN%oWnFtsXK=4!7W0-=2m;8495>vd=ygG-S)?&BP8g3<@2{ zb5p4Vu8H`VtyF*lyVhk!XDAiSc19J^6ryRrQoiQcdpxv}?^Ryt#E{@0}tq!y-Y(vU$L;Ya*ah9YJ?+E3Fe*mo1{W zMuSOWtvy}e0Ubl zd5O;A=9np)scJGymk&H$!E{))CAH*4a?@-EVmG}khBFYqXK&&F20NQ zD_dR#NE8tAmxWSOht%P9htULBaAR}2)JL=)TGc#1L z>Oa8@Y<$OLE+@fjd!niB`&X=$`&0BQQ5>^av2o1z&f+@~;?i6)W)#LSMq(BOCcSU@W zkkq$)x#noE)i%VkiKr<`R!eC4CMV;32j(%O*v}^)xX-)ZDUA?>_o((3jp-q@P<^D{%={O~|+!28?)3lm-@sUc)ygerW9Jq7Yg;txor zD>oo5{iK@_$~|doP+0Tb(#~Uoj=)L#0UlUmVdRe|>4rVb8(EHr&SKE4yAlm!o-=^4 zpH^8BB$!g(UwHQLtXW<7Qd!42G*K(3WoM*ZJuZppaH8B~)ukd+`EyvcMKZ8O2Wmk~ zOH&2!j}5#P7)7nS!{H^G^yqF=pC+kRUpR zF-Ez`+6mE*fWq>_>~JCs3xxJKri=&k$XVPPe?j|?mZHohI36k#W%j-<3iv`X#wnVL zN;P=k%1``1j^xlzvMNmy;bETvhtV88^4R_{p8-LScMAi}uH|)u?MF~*bW&+m><~b` z{Ys4lpgJ|@QWlwUW+b-oR=i+vQTi2g8SJ27Nu@L%+{Z2174yIJ#PbtvWrx+uXf%hZe?Zf+1b2Z2({_XQ&0hL%b zmi}IXr5IiW4to~{a&-A9yvab)z80wVk_Q8KOz=o~Z>Ez4JFHxP4q88F+lfAVgthjhcTX1fN z(hT!Z?B>za3h7(|oiV(Uc=1FPQrWyGKT&nHwJ|$oe9|EwY#TX(E*C)Q*4)DY&u91% zx~u5O|DEd5{ypr<4&HYLOJnhzuPBJ}(BZ>t4;yq)5P0O>W`Lc6ZL_b#omeLDyczL~ zxOn{ABZch~x_pZx=barNDTYI5CZ$`t@O@(mEeS0D_xJGyUkB%4Xo_h%yuf+%P6=NM zr-OGx*($iP_(llglp(;4=<4po>_O0%?ic79?BF*#5nL8gt=zt_7hzU!Eu{qZUB8=$kMtT*cI4T&Za5 z+STGUc2mRhLxwHalV+6w*LJCW&@Iql@sZP}wD{ep5m(n>8If8ODsOAA6TdZIiQjbj zWMXt^UsCfgsR2yK2Zv$$MO}0FKUQ3_#rHR@B_KauPeKEd!UfnHovY>bL3K(q$Kn)8al-I=lx9)ZZk2eBak(iz z)rpeZefqa;IGO6uc2M1!<)frpyh7lUtpFZ&!!+{~h1R8z}Bf=ztk?omvn~<3;Kx|w&%+OpniOVy!W}J+p zw6#`-TNNA2hHDtcwJ!H_r^nTf)5hSU?lx7KVofehV(aslsB?bNxzN})o<7usWoFyW z#b^Ad@ix%$kjRpOJ=MynKc<@fMCFK z&aurd-eZI7x#{{$jAD}YE@jHZ!X=4AT|Z~f=nI+lon#jb5D z4cgzxKLQei*mKzB`Gc?J^yQ>6FplYcA*;C2q}srFDQ7-*bv8#*A!mylWU2UgJB3j# zwkf*@>DBiE<8qPCCjKyW`BP=% zf5QRSLWg_LVMHoPU zED3nVFX72MdVN;Ue`x4*RY^CnzdvI;^AG{tuh=7_mL?wK6nK%)^+PeQ%iDCUlF!|BRJum-Y z%$1{zC)q@g9rt+q*t=2Pl{IOF7Hhh`7h&hV`ml0KKRg|e+m?ARxQW%rJS89Zy}8vl z(yHm{hi+K^&@bpfaX2-#eX^UKn_*FE*=TIBzzQT=JbIIcX67W!q(%SR`1E^xK8Zff zp6xTU3+859b$(fw-sO8ZhrQ!u` z8U2wYK8~=cEpRY&Ff`DQP3ulk?&suY?q zNbXhm?M#x0HM<&)pZfFd!*Jv;$L7H4cQLE^XhRvPP4O1llO1ijx0D_+!Ms8NsBpuh z!v?#Kr`SMJVP|Bs6KKI+%-uKQqnw`YR*PweF62|g7sGb-1?cop#)6IN?ieNi4gN~ey#*JKW=``QM1em$W~<=2R&vv zLDq4NZz2#sB0bK5`~HiZ?FR((@g@1dQ)B$2=8Br<;M+W*!+m^qm_7$cHb!(W*Uq46 zu_`&_{j*zNk+iZY{@F8~wU-4Ad1PI}5Pi7(#=sO8d>n!v6c9QUHkQ@Qu==$>QnZS;ew_Jnx;m?dWrUJ&a+nnb0%NsCh84rz7MjN=;&&|lRpTZ{2=ZJx(9FLX%(ckxGen!hxgq)X2;DYn#2rgVsEDEV1 zZVNj9oL$xInapB;NlC5>?k^>?7r@;4c6{v4)q;`StQTb#?|Y)|_;NT!>#)LoufsAt z9RY?Q5p|dpE)w0^lLYjnE*l>vo!h|(##iIwQqtg3=ThH2*ueRbQQ8&2>7Edr5H%hM zq%*0Uk;udz!weJ{cDo9+B9VnOh5Sdt1(1Pxw_`WbMIGYZhmc%X2S|EDC8Tc@{EXsX97>dcAPaUFS5?DSIHXr zK@_$E;XmSN7-+9#G``%YPYb*qk0o7r4E3&;dSY!qA94CHrl=b4_;Rsi>L_VT^eJeO zYfiX$K`Bw*l^_Y$VPFyyWG=$B3QHJhx$qW7GSqIy$yQ8#0x`V}A`a(|iM2!aZ-mMD zX&GQW`MY0fl>ouEjmTw;xSTsZ>5fV4!o0?t{}817RCxJo0(vyL5k~Vk5A*FF7yX1d z8&1Q4HQi+B%R=kil;^N8Q=_B4bgr8tggK$gkmo4dWp{x{GlcX9s!CJOH-J)b3w`YsqL(oh{`@imSh! zGy;Cmq=U>AP9>ovZdEsU^`u)l5uq#W0#=w4tUh_k%ioXe!RzVndZf6k6otPnzSe&Gx(sDGMBRhFa!qzQAoX`s? z`-eM8o&eCr86&3YXN7vZSh~)*bNi!ziH}GhNAG%Y+W6EuW9<)nhzodOs0)WzCh>D3 zq=^Vd4x4+5fvoum+o|hQY03f@d+d!~*>X|lY;D{s(Y64$ebXqvrTNFifLL z$Tfs)H$RS}ku?&>KpZ{1{%k|E_q>YSjd~*_8*rO2ox#Ogvc=!0<;YOSvy}i(;GUBa zIe5L9d_1<2QM=d}X|5ymC3SfJIHFqOkz~Z$A59NXjG!S~oq*(I&x~P9bzl7Ua)?lJ z0Xv?+4dvG3hOaIaE%-rRKYfbj#1dZPgEG9h7Mm?yuv=53wOb2EH{X%(*{hLmgccv@ z3}Av^SkfiZpmiqyy&|t$=Ea&oaUjV2bF&nHlpKJQw36U*Bmqw{JV(DSwE2sTaqSBB z^|A4Ks|%$v_%HzX22V1fgKNCU+n-3l%fu_8_o(jEP8n?RdFkka*X_%rljx|wTw-ly zx%PzO6)s<}*t~1*y#`AKX~(UKM*6%51JEXkTC&i@vg)b;|EVNyuE5d#mu?Y>r z`H_p69SJjSGgU~+T5q*Me|wO%(1n`BFCT?KAlPn^v!?YAukjUG&9i=?J`szN@6DQu z^SmUMkA<6w^V)JkhP#AnANsqBLJ(je>*|yHP+8?85UVP>R>um^*g%v5^tmk?v^!?IXkc#XTpMx04|z zG;BBFt>05m%qu0R*{THQ)#SZ;N=Sj?pmH~1ep+U7tE0%7zjehD;K#aSXjOIdvRY{n zqL8M*H)d{QVHaE+1LGls-#H~C3FHr{jD?RmrKuIk$C>fQ|FtxJNnxFy4^@vJ_@I+J zBEO60RB2;T;XAjxr%-odM_giy>6_)f&MK4Dpz3w}7{CpW`C)Ak_*6vOMXVKng@yvu zH44(Ill7P@L>_t#ySu7PYpWaZcZG?eZqE^u-6hS&YfQYC*pni%w~nIG1NleF+`7v$ zfUgX4mCp5?;Qpr*aWM-ocN{6f^-d7&|mOBFhl zg!uA!4E95D7Qsk3hg9-nXB1CHrtJc~;d7m)wKI^T%vC}kgUT|@>ez^TNSv^ zE{U);8JSNRN;J^EqUpRinrVOMgSln>U z{wu_at#taPO-0_a7e1(R@=zTHy^Yoq(^mg6z}SKMaG(LJne6h#FSTz1%F>xo9+M*> zG)p}-goIQTkgVUR`!^?fOG%%qb0Ip~zA9;J*8pn;UK9npyoRSAcA}Q9_|FEWNZK{a z$jE2^JyD6B;G&Fbll@u4@HXc7awFaQyn&P#hmZ)B+6#0eIIIj`Rc@p)TOvSI%cd(g zz3JZuVy4l@M)I^f*;M)u7-RB*gGhwwq~5L1>+0cBpuHZf@U2J03dDYirB-gL_a{r$ zttchro7}*H&7$HmTo_veQWnZ+H1cSr1Q~ryYg2uB@r|SLF{YdQ^qfUnw_;W>rsYXw zG&fforGcwe2^Ie_NEcFIM9`K%0Zyj)XK5byZ20$qv;s)kJotLk4mdM~rWi~~04l?Q znBc)T5bP9$Av2&4%0w+G(o%%buC@xli<6fnRAf50A6eo~OqnWo5ZVLMSQz$%Im#P> zI31eI2XYMpt;NEd`OS$~1PEsFhkxf`9Eal{r0{df0P}_8lnfoFYpv*(%A2zmm!^y# zd2Z#IiM=&CG`c=XJn@Z;(T^v3FL28hB8dn<=Fpt;KA9*ztPT}IRTXCzmb5>?ax^;zIWMm0_bmfTu)ml`nL>Ok7nJ2q2gYM)th}-hy_R?ZiqwjmU?P33FpQ+v zadicQ-ScR2Pp5O#j*SEWUevnd|Gp}I(S~;({8mR+>}E9m;s=uGoq{o~*29t1(_ToN zR$ptfjn#!m^e-#b)?VBaeZEJyBZwM6{pW4)jH%f4^K>%dUwf4TYa<9)rL!;ht8Uz@ z@EXU&LoDft_M516_!VD58L#pk-4v|%R`gZa8v01&aX-v`?OB&V3Vq4DJQR2cMkhETlI+iNS|A{9wga!j= zrR|m2ObYXRj1#m=x~oT+6UP9J_~D5D94vg=860NOcVg-wNvNL&#a}a&S@g`yF7&kR zo~6c_&`Bt$zJJDBZQ+m1446%t8-FWW=Op6`uV2y~cWYZgc<`){(_^0hqHh?Bb3Yt* zVVL9)CzvF-*oJgZ*AEI-tScvBw5E>h95z)ueH|M9yXjw_tvZ1{c4yu9(i|@!C%G-T$T(dpn)Al( ze7Sq)+HauDq%yCPBT6E`b`SQ;4uWReHS$H8D`lCE5aQns6)zSQr=i8HZ*7UBFh8uB zav(srKvD({m6#>TS+%pu4bY~BU+wy$9hF-+qSE}A=~UFgRm(D=1Ij~f{Y^+h-08lw zkufI@;z%EgKBrBJ%NHz8-uP6jEsdY;w1%=Y<|ODfbj_Q5Nz>1m7#!^%-(M{adc<2=PYohkgY1pd?a%<>agtK)SYq;E`g6%8e}tHG@i5u) zC>-;*GILUfg8%rtd);0E`-r%LubVM~;Ds78O4gUV@~aKePdW!ic0cC5K}}2EMB?n^ zZZL87rX`t%knEt-u^$c(;dHJ<(n8U1e(FSl1N1T;?+Wf*i!M*m6#Vr&LWY_|Y=wt0 z8cDzf%1mtnxlogQ8e=|+8dYmJ$bO{V04M|TTZA*^n!T5M7!+s#8kyDt%RrJRg96wt zOUCdn7)uE96g7pwOk89$4l3tl7Ba_T*DBxpPH(Ck(JgJ~_eyaqXNFPsT!t`9x3b7p zNhe;Mr)dg19f@!fvAhPPUlJS&H4YY4>?EKWYDV>v6>Ak^664V_BT{Iq>JJZ_)}Kge zbyA7IyP$VE(cL6%&5BpPjdU+$LtdLqdmv1%NuEqs7mH8E`7;$R@eN0MEP5HsbUNh> z=`w9$%(>se!xoo9H#Uj35Sj%u-8_j>i33q-8zjgFilBOcaLqRXwOe_yqdhR8*$_}d zF^GLKbHC>!E^jWrtkf*lSHY}#w+ABj2*4e<<%kMFq>qF0xE}E&rJi7>EZ6bunu+Pe%hK1U%R6bqUtPiO;N04IS_Mu%3rD(# zBOF(p2FM|vJ|udK2S@FT4R2U|Wt)EaO^%2^CzJ_t@2ahcG&?RcD&4vFH*JwwFz}Qt zCaI%+r!`L_=PJ81Z_G!Y+0slS(+i+*+Ybvp_RhU!#Wj`9H@AXAUJiw;BLGt>D2Gb(>(*X_cO+5Yx>)d|6ubGF?AKylnaTvz8f<99c0Ig+0x9{jj zjyT6>=-c@nw>iQCeC+^%c&!}PT6-Ka_o-_P8`_X7DCo|4$U`RG$4RY&31==`OeEl0 znWmx`&Ft;)6W3Ojz~rmF7=zmqBSGVbO$g&-22Nb}6Eum7X}g*I0$}sO%AEJQ4*199O$;{k59 zf=Hcp^zN2^+nIFy;q0fRtt$!957GLNT;CGaM}=$45viu%1@wmUNxsCgI3eU&K)lO= zrqnpl6Wza&>{PbcquG@$|FQ{ASg51{HNQC)P$$S$F3M7j05zi_a9)=aaAPk-r^{1T z`oe2XtX5ey4t>2ir@@_P%n*7Gb6_Zl0`ykbfj$%Jm(GgGn@_^al$9IYB@8>Lp^N(f zQ6y4m_=u)ZV20auq*7kKDowaNK}eocf@32EVnCrZ5<@q9#lvi#-P+Wp=3=4+8=<@8 zcPqV;hd%eODTW+0m{=tNuwxVacDvkpS)0g&BDBqFubQ5~Eu$MI2+kU9y}q{HBk^W%Va_4h%<}M$_bx?ijIBkZ+aY z)HaY-;ApXx06`&_cd*4OTe8ep-R++GSD@qO8ndF{Ry`bAiF2?Ohx?;M0C-83;dyK`Ol&nVOnlVN@facANC$|R2B2Rdc4qls2zE2%Q}BLYsGpPPCfhPr?#_x%m&VLSL*Bc-0obbVzIRr zz@nc*FI?Q4z8_RG#kI-I_e8o`La-O@jdqovxykmIC7ByKxF!qmpRL1Ut(N?eBZVus z6i=p6u7~&1yN3Zm#@YLw4%XuXtYct_u!l&;+0rfMu~``(B!QBUAV=wC!r z;8mLI96mOW(ZvIftaZ1-buN66%!LhR3lK->KU3;K0>U8;;@kdM#YTO zb-5H;R8(a`19)((p9PLmk5%e!k=zORzUK@C6Q&V$B_=wRG4*FHJj$K|hN-*_aLof~&Qn=)C*@E_af|MFp9HA8G*L}~O-8b%1_O{@JG) zTErB3Y9r0G-DWSn2G86b(WL%<{=mk+7+tMRx*2eJVbJ&UGiD6sD=T-g{28i1x4|SN zt*t9V-hO>uuTNn0_no9Pz2@F8FYaUT;u4KL>&(T|FQ@_#tsNS=H_F>0Y8mfipk*Jw zzU8hEOFCy&Vs?@W;zGIx?d6eq!}jjIv+FU)Y3?)ge&YJf2i^sJxqDoi-<6XP0j~c| zwe^$aZG(l7!fwB#_{&lbsw!Y2=BPle7elgy>Mih93+pCr5Z+#L*BiN?3O=qbXSyE? z4b{v`KeS&Zd8=?N=AXM-nPNGsh4l6|Y#E`&|HGg4&UPG}!?)o+X2NNZKLD?PFX7+W z;cwOi#?wUxDY2gbem>VWe1y}wrnz6f1s2VAEXgDv@uwnNj)f0q^{RW%Z&Yij&7kXN z^_m&+>y&qkDMYveOX6NW4|mY=k`WAwMXH9_zMypJ%xJszN&dO(e#Ze6=kp~G(_T8c zyNZ+!qWa9$A>0%A`!ug;yy3WC!WR6_-s5^Cqcx)?|CM@4`Bi=rF!9}?lqAb>Z|ES4 zkGzZ=dPT=f1CxNN0E!)&0{&M-qEJYI{3nFI*gmypU2-nLZ-y=s|89Ku)43ss1SPlUt~T8d~~?vE^z*AZx>k&qApYy{Esa`7>mZmmvj z)Szl_xlFHmSJ34AFAt?E!I^%Skufb(vwHTE&lAqm&l4TIKPzd#o)cGGcH7m~)!|oQ z`U?l5$7Oh_vDTv=N1Cl=rnJ`gpT?=CMkc+<&9-Lfpk{BLZWKVSwa%kld$qr5vPH`B zPugH1nKi7(%ka;wZA`CW%95zR5)zU6r(x4B&su%H>%ObvH#d{|tFRD7y~f19>UM|^ zr~;Y9roJQHo?{V#O$DWdB^cJYQ3pqm*idJ*V~gLFL<)p)2^& zYA%4M7DrKXhS`d?01?!SKs3*aCd=R#(o#X83@jT91r(paW{?+eqddpS9~Le3IQ;%U z?wEJL(*EpZ01d-Z%AMh+#CgrVLP_C%7#jCGF)tyT+c_t<&^RotnQua{NcOwujdl-K z1IXP4^e1&*SZ;HZFsQSPhL+M)%79=Wf4ryP*a^|WAN;84VM*_$<+v_4yl8#i)dymk zV&&p5V)`>z>V)yp*oyIQB&+0PV)Enx`Ho0nN-proj2`cR^l!c&Mah|K0sjz}_7Tcp zMs(n!J8tUEiB8ZwI*!B1ZcG=>-Q+_R(jDmR&>RHw<)h}ne>U)3vFY0H1jj&IfN(3={^cv8iLh)3LyekAq=EytIMRXF}qzD`&}mB zikiv{WHTYHfsYN?S+nOGY0Gp>i1V8|&bt+0uq*Y$@iX)0dxr)2V_40S8Se}`=$yBLG`=LTN! zJD5<~wY%?=?5Y-2ItZ*8%dm?xf7AkWw6^`&R{T&)v!j}Hr-P&28t-1Ff_T>?sZHI| zXt$Z6E%lOl?j#?+#s*3X5>K^WC|Pc<(nvioXqsGXnT^SfdUkbHgQ7;d&^?W)H+T z$(DzQ$;pR@m&w`LmX_I>dcpx}#tycj!76#%#}VzKhbY*niT?h`3dmvge-%q1LqiuK z7-bRHKU1s3)aw}N8#O4h=HoD0TxMeY#_%6SIL{{{i+t~+zp6=Ujx+ZXF6*n|Ha z0N1P(6N|7)i{HNfyioEEr$s6?g`3-Xv%pfE>;3S9RHnhixFnA z-VBq7voJ+9V)Pz98tf44a{9cDIM-B2E0K4{hSMj38HSH=4l48Sxls3$mBi4wT8bi; z;f%&)jsS(HkiLn%$6j;Up&+e(w^_MDmUVgcOwN}pV;8M}XKPD$=A{0`?v=zJe4N+c zyWBU|x7lghU>*UwNCV(Q(Kpj#44&JryDof~lYGK+?pxjY%_y3!$b6@MYTwe6QJdjk zDO}awV=083J5wxLalyg(qqdQ2;fy8HQuf+pe6rNb@WGIeEr~|`yl{{ot-f7ltMigZ zmrFHIK!WXH8h7BYufYfn^iFb~jI|SbtcdhnP*qRW^^|6*+Z+HvAL$0BT48}3A`HjjJS!RGIe69nM89?WSxY)s0^xB;vBte{fg={t3by>1cfBP>?ULO*wC(fWx*e*_HG}V&=sWTPrBq(?C&nU1^I){PN zb!4G?x=l<%Nm=0SkXV+7o^?djLSG&Jo`r6)uZT0rN})41Q?I?hBz^X>mCr931W$IP z6+HVh@K5qTAA&OG+j~7T_zx6Ue0B~MDaUA0DJh6hHu(--B^0~ZnG{#LMp1tZo&t5h z!a`X`ffY88>A6f}gp!HHH53Iu^j4*KKKPmEcIBJB4mDt{h@=Z^y4#cwcFED)(}Iie z^rjiMO8Cb+@;^yWo-&M?k;ON=xm!LLQ8MQX5k32Lf44q%t4tGM$qsQWP#?)%Y>S2u zCV0+79STn;x-GX`HM(G>vrq61>gY|>1(@UBxm6((u{|Tdv!y3QsK%jT&lEeAbc3hs+02pessrzYaSE2BBa1gyaZ za<5a;Q(*m}sz(+)3ZW3Z?j8_wvDlq?zkb%;{|O)oGaXf$z`kS#9{X_+6zv-BdSoGP zkLoB|RXqM39K%R^2fA$Y*q3Xv;wiI&?ZOg_BBJf;<^xsMeh(#0R|z~DtM(sIRc{_a zQoG*ITRfb5STg9-z3h(c$vqh6ZVq&55)>vd!VDV4*^k%6t=OAWIwUA&1&H+bYO~a^ zH2^!;Jc)-DkBG&WkY0%MJd>Dg?yY&RCjl|^G=m}W63YQyb)9de@v#u%_P#sGB9qKr z$|dpfsbpjnSSZua5`Ltt4z4kSBZK%yMQ*0~i;nKr&mp}O%hs*)Yt>d?tQ_Ty)rVGO zkf-Tt4T1)YO~Ub3HI)DGWxi3D%y)Ek34pPg82*9#9kkmOR+N6**r?Tyr15*czQEko z=X{XwdQka*j&I6{I!&S@mEz&;-^QCy&_2wO!t-q2wlh@jLyOx?;r7z+=TJ39+yY)P zo^T37K4(VI3zV+W3cpw9!0efnJPw>Z;`{Iwe*cQ3UPFV^6$VM%V}vOp3l#ptOpYKs_X8ijs?B;=rR^jnDA6xXl5h4P-hvA_9gKNVU&8j z8~W5PIFM%KcExK5ODQ#lsGrN=kSD_X^}_oJ>A!h1w++T!rqv#Y2MVRn4ahXXy3%E3 zqV~ii#34M^701u(l|+b-7&@moNJ4ELgO(KNONE6wH@IL zG$vBV*At|MFInN(M}lKx8{r4wZo4SO9LC@*dYM&%kAapMd5L-H6dSbsqixB3D_(o5 zS8Dv8id<#m>I`!vFOK879ue)a9nEGU#?*NIF}5*U%h5iW(&Z{`PWBB{VnAPpAD4Ii z>*e;%Nyn4y!^qXt(^`r;TVwYD9^EOh0r>jxcQ6r~mikHsg!23?aT1Ay5zLt8-R@#{xqxxkdeTitiJ?38vQ;dtK#3CHi*<*6a1 zrpug)JG99uLTf~nX@CBtE+mrsoK;e5{!X(O%e+Zg?{{q1m*pAw1jYkXPG|Iv;x{

MATbBqi@P5b+U@G z)4XQ%09#4hONi572k2Bz>~ox*hV1ObhLQOT35dB#*mG@_@tN>nAXeJ^3Pxu_U|!m4 zl+ud2aiWi&_uFu7R8)sRg)6;LmaVy&o~UBe;B-!j`n-B+1QsA%5(^9Dl%Yde z-GWmH`<%*!$;NN;p~_WHv7{(iYFLOk=WD0K*t?yGQNT&Uc+ubRZncAf73B}h*DHC$ zMP`x}cV-4w;3d0^QD77Ng0AIKFQ^S%$zglJiu##M{U=Zs5`qXnJZwV92f-=@z*Lff zQR^YH4@8Qr!EPDU!rI8#oU-=_2H^5^$4xMQlH3h=OX9R@@`@?x4AUb>wYl=F>>I|P92?Y9w%g=zj z@IfpCTDcR}TPDzOlbuDgxPoLMAh{s91iug@8N{i@bY)PjB!)84hEiJZHaQ5k<%98=4y(x%tj3sesc?(kNiuSTBjs`yMAsYoT5^g;01Uu6T()EzR z9oX5@w5Rm4cBT%KA!*UFA6RrkB2VDzEG!8NSXm%J*7NZdRd`Rz;P5U)fGtRDLXn0} zkU7??Q}i_>~&hxMcllTeM$ zA6^Z7VuGp;h&hu~)P_cx=%N}KFcuAT(1?p<%D2n)r_UCxr-%Q32j8^s?Gp+05g4Mr zw62B^D*hdfD@MAf!dB`<0tl%Qk#VUUw4gR7HTkMER7AwuV?})37aF&9%O3%TTPhJRQc^xc~!A63tI@OAS#H86gBI7PC*BZ9KZB2(=$Ggc`z ziW<8y+Ar(8JUT^lgX%^iH}5cPQQKWT+OKx-9r2M$3+}@5+LoBvVkH=uX~oY@DTkfd z*0TLC=T$`ulQEB-fRARAhKv~k`l!cYR~8;Xu7uDaAG29s5!m}T`fHEqKqZ^#>D!1IOOf%e!;(Ib7+GzB zG1SQwA7HEYPhVY*ho)cryA2#*nXpS_z|$$AMWma%d923)GEn*AFKM$>1(FT%T+vU^ zib45B{jb;mz^j|D`UKfnzmiti9#wp*&(ldvLt2X9kx^H&;nU_6T4zv-1Rj>Oj&;;t zN{gtH(qHgG&k}@FNvY{1&8gIde!>zB`q@*Hnavr79Y-m5)RP587rBO}l1jC>7?{~r z3#xcd9j1~1)eY}7)X&I6(x^eZLH?r8U>KGho zbDU@bXcPj2c4}*D;LV{VaR-O2g!NW!N)Mj;Gmx+%xZ!6%^9t(aW(USejc;gon_KES z;8bw{X;l>pdixBGgTQkd@1fxYzH=$;}Sx$j%vGf*{D zi9)D*9Vvu{L?JR9x)OzOfJJwSIyjtAmX3NTB_QgEFabg1Xj=il1BGBrY6GzYjmR+r z1>BId01e#;u!NF$1x=OnRal;MY^P*!2xx6F zrV_z3(G^4o$TkiU)~yO&!UoTzn4op57&wL2!O(OP7L4vgtFRa_AT0(PN>g6DSPIf| zX+N>D8)$2=c4(DAK<%P81EShT+X-5qD%egD_Z%$_OAbQ9%>qj7z((<4L-GNd+%T4c zAb%AN%xRnqZi7|=5VpwAma_&|K-MzBj^+4ZJ>4O5((=GwY7ieN?MhD&Wy5J0$SGH2 z1vAE@Mx4B1-s*u1l0}8#xij+K4c##_4H!UwOfqx07~7H&15i4V-G7jbx1 zHLI%IoXJ&vT%5(%=O(?Yx?3NhHdrvcD1RM0m4YYRRalYLhG|&Xr%({}j7WAv9L2^h zNs4moGwTqFP0ZL$vV|s^HQ^Z5HnM^b`{>P)+p?NW!1-VeF;tt zSX>|ad|ZPJPnEW#b~~meSfq&oPH`E=dWhNVm6vwJUSX8eI6~oKgLW~SL5Q6 zvT24>92)cpDPVdusAlJ@Vc8IgRC*&gm}IcevO2 zDWrOj(*l#h2J-_v^)n(Qxqv`@0Q8?NREzgUuk@!R|mWhFo|b(9e&M5s%Jy zCl}))WBK8Da52ve#{E7nLJB>Z&Wm5!@;}XL|2-*ndpWo)+VP_nhof-;$A8PplN2CY zEa(MBgW4-I-O@fgOW1dw=T9f+(;-ZeC-)*?!vTbthEtGuB5^i-hr@>Hu)yAOlN8Nl z+533({^a9I7_m*VsCBZKTav}3WKktqjINQzBxI2_vLN~cS;Cu|@bbt$Bw1SVOUa_^ zWU(E{(l&lqve*r>sP1Gr`hWd!_rp5~b26Pwx62i`#+7*H{6nrFlc>0Y7pzTOp*Ab7 zAVp2Nf(UI#(wk8_+h&N!8>wP%po(wzD=p}yOz}0QaNfnOG!NgED1L(|x;s(M-kd#u z`yyfO-gG?N0!_*{4GB%ueOwCgZ}TYzdsD&_r2OzH^rqHQHYHnpihnQy?|EC<;%;y( z%@8JB;j+#Z-GM8OgLh?0WM_>kraM*sc>3<$tNoQIY>Q5%e4UoC#ocQIY9$N4cu2`2 zDO#cLC&DKY54-`j;+Imz)~Vt)~I8t|QG zwtQPXn#~vcUj{Rr8RjnrWgW^^!XLl>{l%-3pZ8DRa<;(;{R0A2=EWhUM^pXAg0Dxc-^e34)0U-NG{*3^n(DUi1{niVyp$u*-1-D0$NXS@3^-(#+m zqb94#hql@dSAV~-<+`?F84;#Yq!dfr;iIl~+o`(ku0}bGH<}Lf^XcW~ARia=c|IQx z=3nv)dP7{mwd8yGetwX@&hhqlWsOn(ul%pl5^Z*X?CkAwA^R=v%3?h& z6!%IC5zn{n%0|^feMz2{KCcWP7vh-s0e9s#_)<4p&wmOA{t?!y;%L&z^vzIancaT6%L2mfpd^(bEqfmO1}6SDjqXw57V&fsAiEkijZR zR~_EC;XuaFr|NVdV;=HArVYJhy~%o#ExRPyEyH(SYHEF{Io4&VIi4;xC)So4s}`Da ziSQqv9q%53l&62#&1I(V5LK}i4_JFZrF>IBg@1%#$J6Ke0)Oy%M7qM9tgo=k(FEb{ z)ogTGBx@{H=hncf8gP0IIL9mGToRBz`|afU@gAmn+UTLoZJ}d70&eM=lx@f#z#8mU z;or*8zjCp4%*}+GTZ2|}>~7t$`&#amES8r?XZxpbLa5rnH@np+ z9)H$rBHhxfBi&l_>Tb=ee=4K#N&YRtFIDNb@OhTP=lL3ZDuGWY=)K$9dvo@Z60QdD z+_n-|+myH>jxBMEU)`j*c~QqLO-j?WrPAisxXBuBs@=&!t_rsouU|cT`X;!ArnTy~ z=H^%v;}G2Kge6NSl@_1!q|oU`Hy2Vznt!dB6~p}Udb}82jlbvTquKem2#G%>lrh0C z^T~KYw=b{HS6ODL4`#FJ-(}KI)lbvKmwZxOgb$TtQi`vRN9Tj_WV*;dkEhdFI4%ga z*zXhe&)Hx|Z4!&taPAi393H>^?et}Gq7*^hMu@Y1JS^{ZJgiOo@o?>30v2dYEPo`u z(cXBO9f*TZ@9WYT@tjTH*Ekv$Ge+Rq(Yn z)~NN9VkQ13fw3_Ejm%Hcv;@`e~QG*M|^Rb;`S!TzS~M zy9Ka>)4Oi}Ne z>SLy9#KNW)bRmRIQ;`VK5DOv}q7h+XC2VRT!O9~MOFM~8L((QHQujMK-Td;;`R={B z_X5E9cpi{MiJT=15zvblE0bWv6VHT{26g{`#)T zvY2QZ_L48-M9T#)F*#FCl|^y}(N+P zz04=>H;3iH#2;HgJUkPY$9IE+W5Fl?f=|!;nkwSiC3iXTve&&P>d)>3G1TW?AN6JW#At_IdodaN>fS*X)VPomcvN7L%GU{%31(`{>JEFemUO{G`yK`i&(7lrE`QomZ+vQ$) lKpvEbg%0UZlCF*7s@B_%~qMhYTj37G%@ delta 18815 zcmZ6SQ+p*0w`^nEwrxA0j=VKfrznnZ$WEI{=efIJg@P5)hRlG3%OSOnBIKZ8~dWNw+#>+Mx z>72PDs`HOLwzkj-lf;oEw}RGLy>IHeDgvGD zvoV{$j#mvy0i8~LK6RbW6o5Nt?it(m+)&-T^D#Ly>py1RgV~}>-kOITJZT=*g1f&x z(K`6&JGB+`>J7AWGEA+T$-ZS|7ueBSO3S>u%l>+_)Z`?dw}5s#wF1mLxQ)SgHq~v4 zR=_y5I+u0M7>)4q{urK|V#7?plmIWL5kAgqQgYAlA4(agMA337L~}~E(U__f1l0fX zM00#eQYX*LY)+*xc<^KL$EhKIsctp-&AFOwPWr*jZL2!jYk;egq4&O?C?QRjp`X%s z9Hb;OMHb%l$}Q)&PD;aRDZAKlv>9fDu-fADiny~xbi!VjQBqNrA} z218~%ZV;K&YZ(-4aVYgxXI}AlUzq-3m}N?o*g%5A!0g8k4JF% zbTl8=SY)Ht1^lLThWLo>hb=m_p`o1e(o}_5CsdI=}U- z&hQ~~9aVkLX6CqLcmZbgE=az7-1u{SydTuI(GRjDGx5)Q!-F_^lfYN2DzQ-e6=pe% zm6xUW6XSoZb8T<-Qw++6VhsN1=5W;5=HDv@h%aHG?qVF;aq3%kFHS4z7hEA9XOl$* z$8dxG&2jkVXQcV8F!CqV>0%_Z`xt+7iXS1}iuuTn!8xT%RLemyR6c2yW73a51>z`z z|E|K$R~sxK8PO+KEC4aZ2gQ%G#~O-J6!gEo&ASo}^2`CNft1xsLuf}z2n4(#uFm2anO~<49i7(=GX}2z3CtiNKbga%AIPRw^s+`_ zVziP;#m5gOqF+V5QQW@qjN(#k5*(SdKDxp+x+>o`RlA**Dm*dD%^pR zgVu_!-8Qy5eG2UU#o7z7q;MquvHiRKWIi>=$2grr!jBFzExlz1H9c9iI0YA?7ZpEQ zT*5JQo*R`m$G~rE9Vmzr{03N`m#O7Fk`VV|;7h%2=Wlz8$7dzb&0J-x88^)sziqtT zYb?`EI&>;XU?k47RY)Y8XTW+_Ty#f*j+dB3I0pt~9|2lz?%?Zlk7a&q1l3^g zQdZ_lP-1{ODSi3K$cXLG6j34^cPpR|8)E;S;PL?OA}>LF)gh}|;d1W<{1T#Wx!M|# z-cE=0tG;Q4sEMwr#<=AG|jZSI7wk`~xtFOq721Ina>J;k6&1spw?> zVY4Cz7X=_q{ z;2~&-Of_mXN>;W=w7GN*mFzcHStk6}|8QanCL^9%y zC|l6ukk@Wo62nfuryIA_i;LTV1ygZrf+*~%b@U=tR;EgxoW^9eN7Z9M$A2!<*+@!r9w(=ymd0`3OyqPNREbqqCioji?S`4-{6t2Wo~(5GuU^^v9?!(*dG!LM7rY)p8m6o;^BaoR{#wls(vguo#X zwXR$xv@bLZ^SP%f`qa&WcnhieAyf-k0k2sH?cY?X&QOM- zy5+@i-z5THh`%y^_iYIHWL)`(*Y8|~j)GyeX8XOMSP)Uh)YB^i^y%`O1OcV1TuJ3| z?v$gXb!&lCC`%TnT>T1JZl8?XzCib2cW(VT+GlQ;0mMP>akkZ_SxSzY{Ks1IZW}|d zWLAI&;$6|XG21qEjuO`iNG2_N_LWNFG-io=VRL9>ek_)QuI{1v$2(TAh|kJVE*NDi zDNKvOB}za72@6&u3RM3%GS0&N4&R6KU(m3F!!mKRO9g%6{UPK^twaWOTOhzdcdga} zs5-MCTRxe25-|31^?h=p02!?mRu)G{sh|_MbO7$1wc1h~9Lt{V=Zhz!w1+C0T8bQR zQE_$Pe{JdlY1R!0NwU_|71*wGRJtCrp7Ho9`i{Hi~0#Ftm6Icy*u-pTs z7+@1IhY6;Ip)6CG4pq?g#ZZjsk{>?qS5_t~eulD*RfjokOhH?uV*kY989MtFy$7O% zB3gyNFkKI~9MRzPcDggJ1a%PwZ9R@PI4$}x7Fs!qI=xVXIu_8q92QvQwG#jN8PKJ> zAh1KSQNDX&UAi4u1B&)ZKczf0Ek6)Qq^P+vb` zDt3f&|51V`hVEcKUr~iz$*QVK!})>zWW+ot?jc~KjWR*ryAth>5fsPeVAWo7ifRvL zKZw5P90ZLaF#*p8hu^xG2e>Y{JnYD<03CR)^j~oYKJwtjVsw89Ol$2cU*Pgws=FL- z#f^V56K=(^a7d7CpVm39G!Zo5dNMt>HW)|8k7d{yQIox&fid3`Saoi2V*1G|y3Gd4 zqO1R~)l2Lmd+|=|r-K;jMK)3e${9Uc>Wusqo(+Vc<$%T)Ln;Mpmu|J*B%WL|foSsN z`BP=%WB9MTPWr{`EuV23xGF72+<%IziE%WcvvyEAgGvkU0(31k{Z0^tyBZ4~BnNi6 z|8}W)>gV1hx?J7$ZMb=Cv_59`5_$`C>l4-w{7HGvQ256EGCn4@fOZt+4I<7HkXXo# zl>hDMgbdu4on*8!MSdkJ&nvf@Kr4G~RW|99htEy$t?f4MMr{c5?WMwr%-^ns+n-qz z6~};CDo+aLMC|u7k=OsyE=wXAufRnl4*>K^8~6V@O}%%Ld#?dW@$~B7wd01VU)e?X zU`6~zp+jUmPN^fUG#r(7J8#8I(%PnLzAj!!zW@v-MN^e_QgZ&{$v3g-GhhnZEQ*y*C?BMTTFs#PshUMgpA$=s)^ zo*y=S3tvVp+`E=!6>RpX&#A4uzZVaPE(a&|baf-7>6h8-s(BWEim*SPk}76l&$cz} zN(7##ztT8jU^JDjEo^O`v^sPs9i`sJVBO~U710*^tBbB>;N{$97aO!9_Skp%5F@h6 z2bH$yH7E|bKtMmO>M1Wkl6UZFU|{y7?$`a6PeI@jg%O~}RpftWv)QSJ<9R`DLn>)$ zeAkd~W7yHi&%0HO-t6VK!TYN47hi!qSp-V3x$Up`yes@+N^&ofA`>tD?ZJbgv^0xP zkG|hMJa0A4rc(f)V<5{V-`iA--?_c!e9S~`zBYXTxsw(x$d{IMV8g9;bA7DUElG^1 zL$**M3Ssu{Xmx-?Y~Xk}!0rRCP%{_@Y2T^meSZk(4llnF_W8|S!*Pp8PUa4q^bdFw zk(!CV2~EZUL6*279(r*AYD%s5_MTo2ZC_<86HKtthk0ncsPnX24z2iBI8u8F8ld!% zuzC){<;xkXgYhpUeZ2jJS0reW>{nQwIAmu3@yfJu1$mrbqsE^|_Fr}&n^jcz7K9ro z-Mm+7p(IiZlV`DeP_6;dt$$s0Ib0 zW9Bek_8J%OQ^#o~5r0w=nSpPHZ@iBq&3ed<_&|1UPWJx~l2h>8G5gLn-5hcOjdA`_ z)U{12jWsT=#W|(Ji3D_jUj`Hz$DbcP5rkHOb_$YfH0f%&xWT(1|LXR22QJ5>6d=Br zWtwGF`E@_J3HOFJY`y%^4Sp}*6IW>K%4g<%3)=scD<$(Q_R$E169o3=L992Q4rL3x z|GvCmZeMObU&c06~qaY4FH!3P$jA*oh7}w#~E0 zWWh^5x0>AlMwQY!-2?WItGDd`Gt4Ki3q(W)L}k!ZaE#`G>+|%#evB(4k;~yU^(@TGYjbwyrnPCjPjcVy z6B;CHa_8x5@vEA2^jTn&*k;|!Vf%SH2J#;=^?=_?1U$^NLjqZ|&YDkD*hr}J>QW!P zNzU~*CKsm>8Qy71Zx(B@R%WMXr*m@lJTlm>JfNA4kbC~6Xf}0uqgtMw{<2#Gz~j~b z{cACOdN!TIsE)VyVo_!iPHDgMRC5<#jKvFI6seV{`IN^BX-?4m=OcMqk{o(GA%>X#rl38~!P)5S2P(EC0Hhxpbjc!65K>TIdz*0LnR9-r8`)5+~v7E|~ z+=#?j=SuoQ*dweuedt(`Th8(Od3ZLxYSur2R&RDm8i1}^iXfSv78;FAs{TfbSLbt6 z(ed+l_K-*|N@rVcHFI&-LsNLGDu zsC$;eKu$xJo<*NSa8kCT_k#eBXrpxez5W9S0H%TZImxCC)pCgOqg51FQSUx?0LiYB zjh!-snT_WEqNy?0vR9AMT%645S9|*~Q;Wq_M_3lwT$Gd?uczlPJ@+;NF|9M;!l;Kx zNsaCG{jGQWnUUjUB~=9Ht=$_j!II12dED+c?CJ{^N1IX092- z#pm#t=WFbOf~A|jKp;w}Uy9%=kCu`TKpEo4f6 z>K6v|@uF$d7MR@%XgXRl-1!4{_rWyr*)^K!2zzweLtg0)c(9^;u0O*Yu@&c{r`9aP zf6eSh)T`zt{65b>nu`0d$nJ1}t3XPAI4T9)Dc8w(1hOv|0AF+^<{xAyD;AwlfE(sx zzDw0ilw4xad(QoIaS)6q(#y>ycD*%27zfUAkBQMO%YhKyi&Fv8B ze*p%?qzGOB91^BZs{xDVYpyES1%pHy)O_e1B)x5D>G^K!%~0CXC>NKl1d60TGX8XI zLGc)SVm9sp7Q$MknWaGd#8(*Fjz@3b&1awz-hRlPj(a2QTK@n$2Rbw^WxBV2L)bf) zszu)YkU-4Jp)SgxT-IM29H`91rf|g%s3;>d&gRs2V=Y@zloyv`PXK1-?e1cDi}cO%<-@Qq-A zT9t$3UXD#zCwgoKKChsV%0Qeh;X)osF%8=olP-W#GS6BPJE_2HztKczG~zF_w`ZRd z)_e?(fMuVO|CVxxjef*Yx{DU^AH`}EOX||pbaYm`FH}lzD7A#kQShI$M6!w037Jc> z-gGln7xb@_aMmP~1?sl(Ld4XI)7;bWpthS^;7OSs^YzR?no|zffAOprhcylEIeeq} zN-Y;@Wo6PtsY)5Ahul50`Sh*yZd!1YK5W26{UzfsAlP}NuYFsAy=!K1a?_@998V{n5XqfwVR{l zV~KcMM3V4^9H}gor`Gj-QZ$3DDb7KWt((6(K%SBsQn9NCJuxLy=|gR3gNw_rZVO;I zl;JfB4wF^-4Z`BN9MbFdFa^_%}Y;yVTevff--9n5oBUkzjdeyfL0qYs6x_ zu$_rvHtyYB+cYT1u@}|b1JLf}Nzo$R@dqD@8g6UDo!DS?h>Iv+eFe%G7d!C83-UB= zg|Gzh#?$DyF?sQ2#Q_;Q_SPQDL16INwHK@vBpptPQhB#*<0bA5&Vp9#!nOvam4Q34 zXaFQ%;AhEt6jd&f$7WY!(wJ!aFBSC$a)9(5uYw189;;5~SgR%T0cmLJR(=Kw1b=BI zTWuoB!-gK7zf2lkkI+@BeSD18iqLC~9j}Dwbot`Zh?@FvT}>Tq`(L%0he$0q{qIUG zQXu_N0~6F)F2FCIr5DneJNQ4G#wjr7U#wiw;#!HWJS!N$> zCIUt;r!T}6Vh#0Ax@Sedw)dH=7}pTpDCzRq0cMy5)RcNw^{xvlY(i0_YiAs&A})LK zPksE-;WNw{1Q+bjbuI#O)K`PWlKGoLrT*uf?-8I*~RfIzNbKu*&$ zFv8slysH~jOp%rE`9EjFZai5H$p59+$e{*{4|S}D+W=6#DwSni^}12kHK9NE#s`yHA@=~WF|h?4;}-5G6xuR0w-LpQhiW|C3Mq7^>E45 z(g4U;rNb{bRvn=7rAE9X7caPWRi$Jz`m5*h0X2^<={3$J>|eS;G)mGJJC)Lz5rVLj zu~1%P&2|S@D=+PiSmO5vzGCcbBpl#L5f+h^^K;mR-n@UoX#dG696KDhbGelRSByw8 z&RntB^P#j>ZA>L}&w^cYNxfY86^cH{F~e#CF487*K=V7*czP3iM2kUiASA*pPyBp?7w0+ui1GyCk2Y|DR=@$U91=7UE*1^u{WyExY4%z2Yg4`+vqR=moPY)k`|&t zoc)s?$Hg(8Jn8c8o+4)u_|Gb)wR#KJ57k4Q&szduXZqm2_|cjVqUnE)`Ajsx&AEo z)~mLU$`Q|a{=#Jhw{3F-kKEvExNtTqOERlNd~Niz#+(kQJB7H6(YLdw$E$(iR3E`q zyLwoDto~03Q{ah`dH8Qxnq=9P5RqD5L6A>Ii%$}}>r;sjRMR|Qd9}4wLq0FNU3s_@ z?SwDCS5$anO_&RQl~K+!)5!9I_3fAGf=A~PAF`VcJjp>>7gZHPPCAm{dt5y~Gpg}_ zqMu*1wb^I+a5d*pMHo)9){qR$@KVPr!}_L*2cl!6dFI^-vuiS!=e+!2*4PuGK8uqd z{872Ft<$IX`Urq=elk2yat3uL8BWHDA5i()j|0Cw9?1H6AVVZ!>EVweYOO9T>xa)` za3=xgbi5cG4}%0pQ>G2M+>Gv2Yxt^s#Se#&u?9Xl$)H~taiR4l;KJD$@4!YiO?< z8=U7F++NFw&>;8h!zyvWG+R#qe zpclB`htJPI!vG4jzOurTVgi~%>e4bv!GbjO7lj~{t;K7~o~QPwbGh3GdUPnV=IgXo2DnOPEx^yjd*lnfbGD^}p;iwe2 ze$7&BPvZB8YWdxUyYaAYKKng{;i&lrtuuv{=r<|uuLkg7($LHgQy=u;4{mE& zbNrFW-<{yO?@FM&zi&6cD7i8IM%27&I8Ag?wDZX2n-u3~EJUI2wiEWhfsPpK31u6$(%N(s8u zGmxl(+l07PT))|ce;u21-WEO0Cv~GSr=nrYmRCdo^jN?Er(7_wYnZLp`_m1epwTkjPUb}hnEBYQSDLN#5J97s--oa5=MBt)K`PNUN&8FN}bXKBA3_YokaI6OKDTvV-r*ft2aCv zDN9v0pyouYLjOi_DXV`*=mOuP<-kCdh1^w()Qtp6lw!o~Pi!{aaZn6@%m0R!9eL8Ur!0tfz8L&taLy1)S+| zkxnh_c#vG{)ZW0QS{(ge(=^1Ba0g$w!}pL;_X~yDE0*6HvsyGjs*t6r{Hi3KoJ#*? zRJcOPkvlz6jS_2ewe8D;$}CXtL%bf6!jZmk9Kt?eA>HT)7vjk>2^T$=pF*{&d`d~SkFD2oCSeZOp1Tp#x%4F!h71{@{p5TuraI@~`52@@8~J$8*8>X8X+W=Ni80X#@Z3nmjI`o;|*;`$ohpe~m9iI08-! zX|+(U?qvwJc!qH$47b6fa_NuXhp~rA6dhj7X}O{B49z7a2YVo_eM~?YqF)qVCk!>ZhA({5eQ6cKz)vP42_)?YVkGMD`uRA$m&Y0j`1URh! z_R#uD6KMwiZi^zaMH^bIt*mi44Ek1c@nve|dU!9R7eXUD0)prjTG*&Sz|zBt3gv;$ zG8g{CRpBgb6m2x@EgoJZ!BJ-f)F_H5V=+~_B@cA(*pgg}jI;lHsO(H0%duuI0W7OG z3q%Xhm|djzN*%iM@+8;Uu_gKUQ$76P37xH$d);`tc-dG_`G;kirp|ntW(V5vg&8|4 z9142OJy!b!{-&D~uOx~lx?F&ZAe#Q=Xajt!mE-^vb7y^Cn`R|Vx!5291N&znH$|5X zr#>9s0mS@vuCqY+1%SM~b;|3YdsJdoNOKPOr&WlypoPjICSa{s`LcOx)E#Z-?MtqId8>K&Q^tZ7tW zUR_v&msj5RZ*{+Zde%A+X><6mu~7KICvgZH6V>x2x(Fss&@m-UzjSz?>UO@{*A3mC zLLOfmhuYV1ZFoRhw9a;k>8rBvS1_3=a(=g=1xDzeJxRi>wUAg*b2Jx$lFEEK-sJ7$ zpY?iVjt^FZ){NRm{|u&F zWzRyD!$d^q!Y%NBq{&z-mqK}Q>oKnB!Inr_K6>BF+&|E&Y=N`QmTJk!K(7C_k^0Gl z4k0WE;r9Z+;`7R*L?Bzx#B@Z{7S>wn>Q?9x%f<0X;a(Qn3kaKuJOcgxHUe{a0t~Rq z?x+rlZwl!$CAKfehIn@Bz?{Q$D|$&V9Cl-}14>F|ADP zAo-oCIbwUp`55jU(i2N?9IVHJJ%Ktdfq~vgl$76~YAKcv>w6>?H%ll1{F>=yu6i?{ zm_utm0^vPtnxLVeXuxKJ<8>j^o~S}8qyDs<(2M)Jl(2G}pRDR`^(uGhbI@^X$O}Km zWC4s4`j?xVB4;iQ;7Sf4Q}_tYe$=}*jVFNwZGo)Ft5&|H=aLMUXm-z zSRe8>Fo5?%a%2-!0j8q!u z9|k=NH{x*tN55BjG)h@YvBqwoN~%)sgt2=-M&__@!b6j;TrGtM+Mq;-f2@9R?}kvG z6bi*76l;_dSvhKbXy7~iAPRjXGA!{>#yc=mOZ*+xHp+dZB2@Sh7gfZLf(tch1dA+k z17Jt)fl`2e=K+z>KwbZg^*H6ledDo#`|kM2yve`aZ`nWhjSXEjj2!^yN~|uc z?_FYtCP?eb#_64qQjxN8+4O94zON5g2ba<4Q#Hyr9BdQUbEhlCT#QTQ))yt*1`6a| zj>fz`79)XaLmVN#6pIn%>WbVsN+JH84HQNWnvi8(jHI*bc}1ZlKfj zmXRFvyh&kzv%^c$HN%HVLf`EPaTKWCpw!NvqdPy$4}e-@)J@G63uCN^>$Z^?|8$qJ z-`*-&cl^T)0J{R83?apGAhO2CE7hxY(6aTysyvU&KP}nrh=gLaV!M9eIKfNc0`(zI zoL4>xDKE5M{SXD$v2l#jm^%r1qsDJE;%}QC_R|8RuW5M3-Z&J0!|ocM5?B}c+L?FV6 zyg=NntrPQy53qK`&Pv5~vVllQ4Ur$`nc$Z`f)_|EOd4S;JZBas?mJ?b?eA+WZeij0 zqueI_h~Mtk_B;@Qihz+Gul}R71_F&VA5VWbyUnzg0Gwdk@Z)}^S#tCJ!8eTfLCp)A z4Xdelbq}JuVcrgiVN7v8j=xtz=q&?I;SLJsBz_S1MwotNx{Bt$4nV9hkPHw=t=il4 zV*=U9{U_!9jT|>}D~i~ke_<%dYG768c8J+k=bUD-Wb-*uicWT`-!OT3!eTYc46*{3 zb(J=j_7?Re)QUvfdWu4#MZ!_?VDxBu6$Rh%Jt(dLFabI71ig+%*f$M6-@yYZg>|;O z`{r{j<7{DE>FS7vNWj(t$g2v|>GY-szwv0sZ!h?G*ObYQT{XbN2!M}^lo-0vw6T(g zL%6%riBNicQ8p)WSG;8~a=P4w5}U?7sI{W<%*0#WwLq!gx>tf!B_v)@9Lqrd$&_<;21kO@#pd~ zIohn-2~G+wJ0s1W)px|(mqCf`+iTw2Vsa8aRnP@Smv0N3%4V~Gc~lI)j99r`8vto*AFI}rXOMmKmBTKA)dC-`Gtg3+i% zX0TCp8`Bm8U!uk%rd^b#^(N?@Y}gw&eDA+noqDKHp=C;bU$?kWf|cOp0QPSC!jcP9 zSqNLnjFJC%WO28CL44tsZm?a;)RkPxZy5pEvqX^qedQP4k55bVDh*8no((D3f3_66 z{C%fv{J_7D;u7S8JIXf`a7~(G8}+pafEZhA4AEeu74=raivCKgP$DZ?Ym992mK;*B z;xy1GQeU)CXHgU|*0`!$wfG>_A|IqF$t>8`9Y^j?qRz0k-T6k$SA3czN!fr4q0yU= znCi7mm@qwvPBnnPPKvgn=X*^8pU}winJJi}IS@keA*aL@C2Y?@QXwtvZ=n&Kh<0%* z_1p&sGagNl`}L5i6)QhcQ_D7=zku1C`fVi1Hs33}4`GyrCa=~6XPEDYwrKzB<2A$q z1Y0BH3;_g2{!4)?YaAr6peXSaUSUsi>}+*N(rrgH2VVN zEin5z!UrL}gi(9}Zg(>srJ62<{wp>^IR;^3{9?-u9upzqtx1rsWO{LDyPRC)%OJ{9 zqOrp$LH;pThb0zDaA28*6Nr^R<|5kvH3FVAQ3n0!o)(sEqPy_l-r|PAaUy~c#anBM zi4Sym!NS&%U@?%Zpm#U0aIn>y<(BTL6S&2PBqfe^Jw=Y#sWHZUaW?iyER0u)!J zRuUevgN2{0ez$Wr&ZnSFhA)UNF2caOdLfk2ha*%HPKedBEE25I-&yb|t*gz)$2Tuj zp9DmkCrX)XoUdwJ{oz_U<&f_LhfA0mJ04gjyv zx)9t7(38z`?TNWDc3v3gKig87oxJrb_MUKnfR%cmxWgkP3u%=D+pH#9AmBizOK=(j z)TfCe$8aSkFcHPTtJhqz7vf$k3Y=LYVXVnGx;kg{UB3i`PL=-g=H;A~+S|Gr+6 z4@-ld?-t^63%U&ayDYRG*mX?5V#n^<&<*cWTz{G8J}lM}EpfE9^*uD!m)3@03edF@ zmF0|vmqJ*(0Ge1wm(18W*pDU0z4Qb=jfHgB|ImKGHq!Q8lL6?}9MS;tjyfQy=w)e^G<{c^F!RzJyp!FDA7GO!RmHGGRn`^`? z>=OxrpfK2d!2-ic@fMvh41}}_LBg>S2xw@ZUw|Dy%v>R`Knjs33B`d)4p&)9To@s2 zL*9t1VNuG!E`{FrE} zy(_f+t+r)~;cEIg2!KJBPjR8s7cWtJS`hgu0xhi9(~2Z)c+RG*0d7j8T%~mL$Y_M* zi?f>szf}w?-ow?GoL(Ts-AK=TYr6*LYieI8y;*1TeH5mMH$~{x)!LXY3b?d<>~U^u zLB|1#0Mj8Dbd!npwcNLk4hlRNV=(m}A~RdLEOYqe>^eWP=UXOd1tPdqldyw=7-4No zW68u#0~uRH_?^kxKrmLzrF)uc^|GBXHtc9Nts6-TF=b3uD4SM=io>0i>-dPbN(8+d zvwI+Lc!+lKu;8wsh3H#>wKzUWE+>|)Uy+Lyf z!}TgRr-EjZnMN5Vhm0ggRS5fVO9acrHU$YlJJ@@cO7)i;$lJ!xNqx}#N;4TOwD3%n zP}nZ-*RsVIBYCg&z;6|@q1>u=BGew5yf+vjEpU3?(>5)i!OUtoqAPtjh}nx3P;E*T zo-F<*5mupV>8I$MVzpM2v=V|6W!aA>6L3whke9)~A7la=BRG+E5GH^t z`p=qY&+XzKNH8VoFx78=cv;KqcTMJ&vA3?|X}9smB#8!HJ$c1QzWjD{slGZ{gb#y8 z7GMS(#&p;q;+(Fr@62(y3-lQii&hH)$heg1nDGao0!~+9bD_xFp9tigIn~53`<4JX z{JMu)ULSQfLRS*rFNY;kANaHBZS~ngIm;Mj(rIMCt?J2EZB`n37gG49C)b$MZ~9#g zA-c*Ffpo=)>}>ykkIU8WQ?-KK1_IT6P_5O>Mi|Qu_s5g`-1cgBL6@+l+NR1m$nNy` z_VOgFSnA@4ILNJf?`KdW=`i;mning)P{Av;QQXwSbN zz0i90I7=Grs>Lha2(+51xOg5DVe2dow#^Rg`AKVu`;!^DCc|P+AMjoO=N}L$hbE(D zJ6D)bAlv^#y`#F?uK3(&0rS5z+gs_&5bjLrPw5EL>n5HkHx*_cg0RNfxG}0IX~*{r z`whXxk-~G6IFmc@!@;0N7RT4)9E=#W$JnaaNhPTCse0#9XIh-Iu(H$zmieCX6q%0o ziWJAfp0Y4-GYAr(%2u2JB$=aRSEU-k4YD~ws{=4MSDPdR1eO-*G}@8{+*;cCNSXiA)6!wgNbE%9 z^#gzjBjOoE{-Ug|BRgb~2m;oBnK{8SZq*|-d${ymY??&XbS$(pQ!@%qABs zOoI7Zahf<@JjICo84k1**bxf~^++KLDR!oWE=nB@dTgjyT?|fKDM??Kn`>_dpzwotDi(Jz8i|#2p#Ku1Ft@KY__syvjowiO$qfiVId1 zjS&veqHH+Fy^ODgwl~(iZ2Q-c!*p{jwl@l z=^a?yxPVNPmhTpcH7eaqXBD3cgO>`Xpr%C1D66E7VqI=UD+#=kaGaGZhYU+AM_@NwI?8fv3i0Bq9ZA$7 z*VYGsZM$hYI*b4M+n%++V16!!murz_b@Q!;3>ZiB&JLufg5PSZ@+amvblyWh^mgrV zAac<1?@*f^-eru~HhZsrBOt6_)U8GJ7r---PDBl3mWQX`Y-0J%Tq|+x#OT1S_k8CT z!55|(2j4!aTg$k^wzKQv<>`G*q73)rfMYF7B)(pGqGCY_`}J9h6A$IP>?rteRVnf0 z7w_(OEe?>Oyjd#d$d5HIc8(bRB4Yh~^k%XBaD8$0;rX=vW?t$B`}_QOyz$M7j~9m_ zS_=9HLk=f=LRPviYuosS`2EK@mnX5~gbJlsSPNln8UAypIX#pV^}jz%FsDyE@oe?| zQ|GsgujfxBi%#tsj4dcQ+=EuS`gId=8d5;)Rywdh2z+D#%^B4hH2~Hi`uk}62bRxr z(_grS%|FUk$9uAO^`#63r>Mp%B@%YcB7~?#wXVj>4mG-m+fFa0T3-wgE}pH8cjN?$ zqtb;6(o`+-ni8afHWJ}>%O(0_upeP8TD#*nO>oGPMus9l_Qln}%&vx*R$%QBkffpI z&j|$RU{pe|qTYXOmw=boGvsOrDrVcEEW?Rm8=}MTXtO&cya03k|nPA)2Y-xdc z0fCeN4Iqb`;#`!rElyf1rN#~@DI+|yl5{x3bR59`n|dM8>j7<_dSlqc%rKTAx0VrCE(Ib^ z13r!X=6N@0Khqjh52yRo1t~sAhbgPAsK$L1lr@+>6#g(2Dl9M1LiYLp_C@Nh8WKX! zPJ^3j;4v*MNXiVMAe0t!*dErQ!Bj>V6`?51)@l^QXOomHY>dsVfb|j6Qiw!%^9cIQ zN%wu?3snk@9VeQdNQG*dwW?sO{{Wo!CqG)relox%K$kDQ1OeY=Ts^#s^lC`~5lG1I zC6EMz>79zQ-etjObMzypBamNtQzGe<2ZL>hX^wRygdL&{WGPl{6&k zWr;Ut5#&opL*ulw`?_#46m<3Rn3v;Rpf@zF;W3oiI&`x&r1W2^qm0 z@@63#|JMbn8$gkd$b+|lTPG|u1NXV~Eb)!@I_Y!to)27AqJ$}mKduKMb)$euI$}vB zIg!iEFAQ0xHNTXOSb_+Shk;MDBCw|)@C-yrmGKwx2t?(zP=Hn`#6t|n3)T{o=>Vga z4N)7RY$IWlE-TsPN55H;6b6=(hv@x!RvQ#(w3ML?5f^cc;Z1|dkIEO}3$;7O7<5p7 zK;!bF8im-bD;CFlL5eO%|6S-t>;q{!5Nd|Nm0IrS(TFi}ssc1A;~rQW@R$II6nZS? z@bZDhcA=XU&|N5?{vp7PpoAtpv>*lC3`jl}ia&v7@36>;qQpR5aR7N57Z)+Tq&LVc zb?fo;K&_`C1k<0vUsuGgD4qulY*2`4ZUKrY*27iU1GJj1ken*T(h)I8*BZhdBHCs+ z`TXSf@dn<%vuS#qW25SR{`N=R z)!nbKHYt#!F-wiAv_P*WeaM(vF1g_$p|IdWONl0YI?|(bieH_tS?;LO@>z{JR`B{9 zw#lYu+Y^A?B=}=`+%Z%Q`&yf0E6vI6RlYF7hT4BTznQ~S(Q6sEt2sl|!&F{vhU!A` zf-tx&QMu1UCnxkLW8xIos8$>nN>uT>V0uamb9(;zkL>C1l>u{qNo}@+`efBE;mQU0 zd|9I=50rrWloM<9ac5#5u9uv5OA+t?d>)<&0&m*Jv9_wys3+aH#8Y{-0tt&8L$h zU;I7IKaai^dGYV-!8rdk8=U`DEW((w9!CxGPx*O1%!~XY|B{dLaekRk@@amR&+>V` z$glIS`L`UFv;tNtzT2wGikhnAnyQ3uQPsP%-TjyEvDV2^ldj}L>uQIqU)a)`tyso^ zY3wL}#ZqH$Hjb}&&PxLm;8cmg$p>8d@tY65AxSJ zZhTkT80G)U|0*@nriYqry1Cc)2k)PqyxzlPj|VTGpTec>{ods6)$QHA>9jRd+*S*?Dr9I=x_i80yy5PB-Q9nN-zJZL zJx{!Jjiqj;cW`j@^uvc`roTngn3Hu2yBtjr>{g?AXs-!$ORtV3Yt5^FyEU)= zsf@-a`L{%VDNDB%pJ%D~JYN%^O2nrV>Ald=me3 ztR_0$=;lJ|NUIgIVwhiEj~AnVtMT{zd^9^B7a`rJgfb@hWj+}nsQL2xe3fN}`d~Jj z{#_>hRQ)tve90%pMR++mCZ%}OcyvA(Po|6f^LRR)h2w%yi{n0F|C|km)F;tsP0rmy zoWtYSzn#8J&XJ-}w-Ms39}3HR9SUpHekfeKL4bkA#6r?N_Qn-<01jS%xWmhYx?D@0TG50TY*3e*qdoIW#k7Fg0anF*7hSF*jv2GC4S6VK8GiWM*Py zWI1CqK0G-zGiER~Wo9umFfuVWWi&E5IAUQiV>e`GVq|1FV>6ete*s5-Fg`vCb98cL zVQmU{+BMQkY)oMQ#qo2h)zUIriqdLPGqiN5S0A(}T92aBs%b|xM8rZwtY+UgYu9(LLdgHL7qF~E|V!D*5`$C z7E$-XOOY&>!*Ymddhe? z(`4{jYw-CC-_t})zH-+RZzkRAh^bO{9r1SBy`Fei?OsNF{OjIGd>U}q6JN9LEm41S zHxPfC-R)6NxR;ZE#pUiDQJ;0MAS-{md&sK7$yrIp3*X;ICQ96mWLwgGH0m4fRb+RE zJ0101cbptdxie80&P;+F{oy`AX1d&~qw`$m-X^=`UU@+F%foU&4wBg%@E Tensor: + u = t.uop.copy_to_device(self.devices_4)._shard(0, self.rng0)._shard(1, self.rng1).unshard((0, 1), (self.rng0, self.rng1)) + return Tensor(u) + + def test_2d_shard_basic(self): + ref = Tensor.arange(16).reshape(4, 4).contiguous().realize() + t = self._shard_2d(ref) + out = t.contiguous().realize() + np.testing.assert_equal(out.numpy(), ref.numpy()) + + def test_2d_shard_elementwise(self): + ref = Tensor.arange(16).reshape(4, 4).contiguous().realize() + t = self._shard_2d(ref) + out = (t + 1).contiguous().realize() + np.testing.assert_equal(out.numpy(), ref.numpy() + 1) + + def test_2d_shard_sum_all(self): + ref = Tensor.arange(16).reshape(4, 4).contiguous().realize() + t = self._shard_2d(ref) + out = t.sum().contiguous().realize() + np.testing.assert_equal(out.numpy(), np.array(ref.numpy().sum())) + + def test_2d_shard_sum_non_sharded_axis(self): + ref = Tensor.arange(4*4*2).reshape(4, 4, 2).contiguous().realize() + t = self._shard_2d(ref) + out = t.sum(axis=2).contiguous().realize() + np.testing.assert_equal(out.numpy(), ref.numpy().sum(axis=2)) + + def test_2d_shard_matmul(self): + a = Tensor.arange(16).reshape(4, 4).contiguous().realize() + b = Tensor.arange(16).reshape(4, 4).contiguous().realize() + a_s = self._shard_2d(a) + b_s = self._shard_2d(b) + out = (a_s @ b_s).contiguous().realize() + np.testing.assert_equal(out.numpy(), a.numpy() @ b.numpy()) + @unittest.skipIf(not_support_multi_device(), "need multi") class TestMultiTransformer(unittest.TestCase): @needs_second_gpu diff --git a/test/unit/test_multitensor.py b/test/unit/test_multitensor.py index d87214934d..60f4e6701c 100644 --- a/test/unit/test_multitensor.py +++ b/test/unit/test_multitensor.py @@ -597,7 +597,7 @@ class TestShrinkMultiTensorShardedAxis(unittest.TestCase): t = Tensor.arange(64).reshape(8, 8).clone().realize() t.shard_([f"{Device.DEFAULT}:{i}" for i in range(4)], axis=0) - with self.assertRaises(AssertionError): + with self.assertRaises(RuntimeError): # sharded axis shrink on non-device boundry is not allowed a = t.shrink(((0, 3), (0, 8))).contiguous() a.schedule_linear() diff --git a/tinygrad/callify.py b/tinygrad/callify.py index 2877eb119f..2c9345af9b 100644 --- a/tinygrad/callify.py +++ b/tinygrad/callify.py @@ -83,7 +83,7 @@ def contiguous_mops_to_view(c:UOp, src:UOp): resolved = graph_rewrite(src, multi_pm, name="multi_buffer_view") if resolved.op is not Ops.UNSHARD: return None if (view := _make_buffer_view(resolved.src[0])) is None: return None - return view.reshape(resolved.src[0].shape).unshard(resolved.arg, resolved.src[1]).contiguous(tag=c.tag) + return view.reshape(resolved.src[0].shape).unshard(resolved.arg, resolved.src[1:]).contiguous(tag=c.tag) return None diff --git a/tinygrad/schedule/multi.py b/tinygrad/schedule/multi.py index b85ebe4639..32cc82485c 100644 --- a/tinygrad/schedule/multi.py +++ b/tinygrad/schedule/multi.py @@ -1,12 +1,9 @@ from tinygrad.helpers import all_same, prod, getenv, ALLREDUCE_CAST from tinygrad.uop.ops import Ops, UOp, PatternMatcher, UPat, GroupOp, AxisType, graph_rewrite, broadcast_axes, _broadcast_shape, sint_to_uop +from tinygrad.uop.ops import sint, ssimplify from tinygrad.dtype import dtypes from tinygrad.schedule.allreduce import handle_allreduce -def shard_count(multi:UOp) -> int: - # the shard count: the device count for multi-device UNSHARDs, the range size otherwise (e.g. threads for LOCAL fragments) - return len(multi.device) if isinstance(multi.device, tuple) else int(multi.src[1].vmax)+1 - # ***** multi rewrite MSELECT/MSTACK ***** def _apply_shrink(marg, s:UOp, i:int) -> UOp: @@ -75,6 +72,13 @@ def shard_srcs(msrcs:tuple[UOp, ...], axis:int) -> list[UOp]: return srcs def alu_multi(root:UOp): + multis = [m for m in root.src if m.op is Ops.UNSHARD] + if not multis: return None + sharding = multis[0].sharding + if len(multis) == len(root.src) and all(m.sharding == sharding for m in multis): + srcs = [m.src[0] for m in root.src] + return srcs[0].alu(root.op, *srcs[1:]).unshard(multis[0].arg, multis[0].src[1:]) + # resharding: single-axis fallback via shard_srcs axis = root.axis assert axis is not None srcs = shard_srcs(root.src, axis) @@ -82,86 +86,145 @@ def alu_multi(root:UOp): def reduce_multi(root:UOp, multi:UOp): op, num_axes = root.arg - if multi.axis is not None and multi.axis < num_axes: - local = multi.src[0]._rop(op, tuple(range(num_axes))) - # allreduce in pre-cast dtype when sum_acc_dtype promoted from bf16/half + sharding = multi.sharding + reduced = [(ax, rng) for ax, rng in sharding if ax < num_axes] + remaining = [(ax, rng) for ax, rng in sharding if ax >= num_axes] + local = multi.src[0]._rop(op, tuple(range(num_axes))) + if reduced: + assert not remaining, f"partial allreduce not supported for multi-axis sharding {sharding}" + # all sharded axes are reduced: full allreduce if ALLREDUCE_CAST and multi.src[0].op is Ops.CAST and multi.src[0].src[0].dtype in (dtypes.bfloat16, dtypes.half): orig_dtype = multi.src[0].src[0].dtype return local.cast(orig_dtype).allreduce(op, multi.device).cast(local.dtype) return local.allreduce(op, multi.device) - # reduce on non sharded axes, piecewise is fine. if axis is None this is also correct - new_axis = multi.axis - num_axes if multi.axis is not None else None - return multi.src[0]._rop(op, tuple(range(num_axes))).unshard(new_axis, multi.src[1]) + # no sharded axes reduced: piecewise, keep all remaining sharding + new_axes = tuple(ax - num_axes for ax, _ in remaining) + new_rngs = tuple(rng for _, rng in remaining) + return local.unshard(new_axes, new_rngs) def reshape_multi(root:UOp, multi:UOp): if prod(multi.shape) != prod(new_shape:=root.marg): raise RuntimeError("reshape must maintain prod(shape)") - if (new_axis:=root.axis) is not None: new_shape = tuple(s//shard_count(multi) if a==new_axis else s for a,s in enumerate(new_shape)) - return multi.src[0].reshape(new_shape).unshard(new_axis, multi.src[1]) + # map every sharded axis through the reshape: the axis boundary must survive intact and stay divisible by its shard count + arg_acc:list[sint] = [1] + for s in new_shape: arg_acc.append(ssimplify(arg_acc[-1]*s)) + new_shardings = [] + for ax, rng in multi.sharding: + count = int(rng.vmax)+1 + target = prod(multi.shape[:ax]) + if target not in arg_acc: raise RuntimeError(f"reshape {multi.shape} -> {new_shape} moved items between shards") + new_ax = len(arg_acc) - arg_acc[::-1].index(target) - 1 + if new_shape[new_ax] % count != 0: raise RuntimeError(f"reshape {multi.shape} -> {new_shape} moved items between shards") + new_shardings.append((new_ax, rng)) + new_axs = {a for a, _ in new_shardings} + new_shape = tuple(s//(int(rng.vmax)+1) if a in new_axs else s for a,s in enumerate(new_shape)) + return multi.src[0].reshape(new_shape).unshard(tuple(a for a,_ in new_shardings), tuple(r for _,r in new_shardings)) def expand_multi(root:UOp, multi:UOp): - new_axis = None if multi.axis is None else multi.axis + len(root.marg) - return multi.src[0]._mop(Ops.EXPAND, arg=root.marg).unshard(new_axis, multi.src[1]) + shift = len(root.marg) + return multi.src[0]._mop(Ops.EXPAND, arg=root.marg) \ + .unshard(tuple(ax+shift for ax,_ in multi.sharding), tuple(r for _,r in multi.sharding)) def pad_multi(root:UOp, multi:UOp): - assert multi.axis is None or root.marg[multi.axis] == (0, multi.shape[multi.axis]), f"padding not supported for {root.marg=}" - local_pad = tuple((0, multi.src[0].shape[multi.axis]) if a == multi.axis else s for a,s in enumerate(root.marg)) - return multi.src[0]._mop(Ops.PAD, local_pad).unshard(multi.axis, multi.src[1]) + for ax, _ in multi.sharding: + assert root.marg[ax] == (0, multi.shape[ax]), f"padding not supported for {root.marg=}" + counts = {a for a,_ in multi.sharding} + local_pad = tuple((0, multi.src[0].shape[a]) if a in counts else s for a,s in enumerate(root.marg)) + return multi.src[0]._mop(Ops.PAD, local_pad).unshard(multi.arg, multi.src[1:]) def permute_multi(root:UOp, multi:UOp): # all permutes supported! - return multi.src[0].permute(root.marg).unshard(root.axis, multi.src[1]) + return multi.src[0].permute(root.marg) \ + .unshard(tuple(root.marg.index(ax) for ax,_ in multi.sharding), tuple(r for _,r in multi.sharding)) def shrink_multi(root:UOp, multi:UOp): - if multi.axis is not None: - shard_sz = multi.src[0].shape[multi.axis] - s, l = root.marg[multi.axis] # SHRINK marg is (start, length) - # shrink to exactly this range's own shard: this resolves the UNSHARD into its per-shard view - # (e.g. a fragment indexed by its LOCAL thread range becomes that thread's REG shard, no copy needed) - if sint_to_uop(l).ssimplify() == shard_sz and (sint_to_uop(s)-multi.src[1]*shard_sz).ssimplify() == 0: - non_shard_shrink = tuple((0, multi.src[0].shape[i]) if i == multi.axis else t for i, t in enumerate(root.marg)) - return multi.src[0]._mop(Ops.SHRINK, non_shard_shrink) - shard_bounds = tuple((s,e-s) for s,e in multi.bounds) if multi.axis is not None else () - assert multi.axis is None or root.marg[multi.axis] == (0, multi.shape[multi.axis]) or root.marg[multi.axis] in shard_bounds, \ - f"shrinking not supported for {root.marg=}" - if multi.axis is not None and root.marg[multi.axis] in shard_bounds and root.marg[multi.axis] != (0, multi.shape[multi.axis]): - # NOTE: shrink on the shard axis is only allowed when result is a single partition, denoted by the new real - # we just copy it to all the devices, no real. this will be optimized out later - non_shard_shrink = tuple((0, multi.src[0].shape[i]) if i == multi.axis else s for i, s in enumerate(root.marg)) - return multi.src[0].copy_to_device(multi.device, arg=shard_bounds.index(root.marg[multi.axis]))._mop(Ops.SHRINK, non_shard_shrink) - local_shrink = tuple((0, multi.src[0].shape[multi.axis]) if a == multi.axis else s for a,s in enumerate(root.marg)) - return multi.src[0]._mop(Ops.SHRINK, local_shrink).unshard(multi.axis, multi.src[1]) + # resolve each sharded axis independently: a shrink to exactly this range's own shard resolves the UNSHARD along + # that axis (e.g. a fragment indexed by its LOCAL thread range becomes that thread's REG shard, no copy needed) + local_marg = list(root.marg) + remaining = list(multi.sharding) + for ax, rng in multi.sharding: + shard_sz = multi.src[0].shape[ax] + s, l = root.marg[ax] # SHRINK marg is (start, length) + if sint_to_uop(l).ssimplify() == shard_sz and (sint_to_uop(s)-rng*shard_sz).ssimplify() == 0: + local_marg[ax] = (0, shard_sz) + remaining.remove((ax, rng)) + continue + part_bounds = tuple((i*shard_sz, shard_sz) for i in range(int(rng.vmax)+1)) + if (s, l) == (0, multi.shape[ax]): local_marg[ax] = (0, shard_sz) # full axis stays sharded, shrink the other axes locally + else: + # NOTE: otherwise a shrink on the shard axis is only allowed on the legacy device path, selecting a single + # partition (which is copied to all the devices and optimized out later) + if len(multi.sharding) != 1 or not isinstance(multi.device, tuple) or (s, l) not in part_bounds: + raise RuntimeError(f"shrinking not supported for {root.marg=}") + non_shard_shrink = tuple((0, shard_sz) if i == ax else t for i, t in enumerate(root.marg)) + return multi.src[0].copy_to_device(multi.device, arg=part_bounds.index((s, l)))._mop(Ops.SHRINK, non_shard_shrink) + val = multi.src[0]._mop(Ops.SHRINK, tuple(local_marg)) + return val if not remaining else val.unshard(tuple(a for a,_ in remaining), tuple(r for _,r in remaining)) def flip_multi(root:UOp, multi:UOp): - assert multi.axis is None or not root.marg[multi.axis], "flipping not supported on sharded axis" - return multi.src[0].flip([i for i,x in enumerate(root.marg) if x]).unshard(multi.axis, multi.src[1]) + for ax, _ in multi.sharding: + if root.marg[ax]: raise RuntimeError(f"flipping not supported on sharded axis {ax}") + return multi.src[0].flip([i for i,x in enumerate(root.marg) if x]).unshard(multi.arg, multi.src[1:]) def stack_multi(root:UOp): # STACK adds a leading axis: srcs are sharded one axis below the output + multis = [m for m in root.src if m.op is Ops.UNSHARD] + if not multis: return None + sharding = multis[0].sharding + if all(m.sharding == sharding for m in multis): + srcs = [m.src[0] if m.op is Ops.UNSHARD else m for m in root.src] + new_sharding = tuple((ax+1, rng) for ax, rng in sharding) + return UOp(Ops.STACK, src=tuple(srcs)).unshard(tuple(a for a,_ in new_sharding), tuple(r for _,r in new_sharding)) + # resharding: single-axis fallback axis = root.axis assert axis is not None return UOp(Ops.STACK, src=tuple(shard_srcs(root.src, axis-1))).unshard(axis, next(m.src[1] for m in root.src if m.op is Ops.UNSHARD)) def index_multi(root:UOp, multi:UOp): - # INDEX on UNSHARD: resolve the sharded axis into this range's own shard (idx - rng*shard_sz) - if multi.axis is None: return None - shard_sz = multi.src[0].shape[multi.axis] - local = (root.src[1+multi.axis] - multi.src[1]*shard_sz).simplify() - # the index along the sharded axis must be provably inside this shard - if local.vmin < 0 or local.vmax >= shard_sz: return None - return multi.src[0].index(*root.src[1:1+multi.axis], local, *root.src[2+multi.axis:]) + # INDEX on UNSHARD: resolve each sharded axis into this range's own shard (idx - rng*shard_sz) + idxs = list(root.src[1:]) + for ax, rng in multi.sharding: + shard_sz = multi.src[0].shape[ax] + local = (idxs[ax] - rng*shard_sz).simplify() + # the index along each sharded axis must be provably inside this shard + if local.vmin < 0 or local.vmax >= shard_sz: return None + idxs[ax] = local + return multi.src[0].index(*idxs) + +def _shard_idx(rng:UOp, dev_idx:int) -> int: + drngs = [r for r in rng.ranges if r.arg[-1] is AxisType.DEVICE] + return 0 if not drngs else int(rng.substitute({drngs[0]: drngs[0].const_like(dev_idx)}).ssimplify()) def copy_multi(multi:UOp, device:str | tuple[str, ...]): - assert multi.axis is not None, "all multi ops have axis" + sharding = multi.sharding if isinstance(device, str): - pieces = [multi.src[0].mselect(i).copy_to_device(device) for i in range(len(multi.device))] - return pieces[0].cat(*pieces[1:], dim=multi.axis) - return multi.src[0]._unshard(multi.axis).allreduce(Ops.ADD, device) + # reconstruct by concatenating along each axis from last to first + piece_info: list[tuple[tuple, UOp]] = [] + for i in range(len(multi.device)): + idxs = tuple(_shard_idx(r, i) for _, r in sharding) + piece_info.append((idxs, multi.src[0].mselect(i).copy_to_device(device))) + for j in range(len(sharding) - 1, -1, -1): + ax, rng = sharding[j] + groups: dict[tuple, list[tuple[int, UOp]]] = {} + for idxs, p in piece_info: + key = idxs[:j] + idxs[j+1:] + groups.setdefault(key, []).append((idxs[j], p)) + piece_info = [] + for key in sorted(groups): + grp = sorted(groups[key], key=lambda x: x[0]) + piece_info.append((key, grp[0][1].cat(*[x[1] for x in grp[1:]], dim=ax))) + return piece_info[0][1] + # multi-device target: unshard all axes and allreduce + val = multi.src[0] + for ax, rng in sharding: + bsz = val.shape[ax] + val = val.pad(tuple((0,0) if a != ax else (bsz*rng, bsz*int(rng.vmax) - bsz*rng) for a in range(len(val.shape)))) + return val.allreduce(Ops.ADD, device) -def store_after_multi(dest:UOp, src:UOp): return dest.after(dest.store(src.src[0])).unshard(src.axis, src.src[1]) +def store_after_multi(dest:UOp, src:UOp): return dest.after(dest.store(src.src[0])).unshard(src.arg, src.src[1:]) def passthrough_multi(root:UOp, multi:UOp): new_src = (multi.src[0],)+tuple(x.src[0] if x.op is Ops.UNSHARD else x for x in root.src[1:]) - return UOp(root.op, root.dtype, src=new_src, arg=root.arg).unshard(multi.axis, multi.src[1]) + return UOp(root.op, root.dtype, src=new_src, arg=root.arg).unshard(multi.arg, multi.src[1:]) def rewrite_into_function(call:UOp): if call.arg.precompile: return None @@ -171,7 +234,7 @@ def rewrite_into_function(call:UOp): assert new_body.op is Ops.TUPLE if any(s.op is Ops.UNSHARD for s in new_body.src): shard_call = call.replace(src=(UOp.maketuple(*[s.src[0] if s.op is Ops.UNSHARD else s for s in new_body.src]),)+new_args) - return UOp.maketuple(*[shard_call.gettuple(i).unshard(s.axis, s.src[1]) if s.op is Ops.UNSHARD else shard_call.gettuple(i) + return UOp.maketuple(*[shard_call.gettuple(i).unshard(s.arg, s.src[1:]) if s.op is Ops.UNSHARD else shard_call.gettuple(i) for i, s in enumerate(new_body.src)]) return call.replace(src=(new_body,)+new_args) @@ -195,14 +258,13 @@ multi_pm = PatternMatcher([ (UPat(Ops.AFTER, src=(UPat(Ops.UNSHARD), UPat(Ops.STORE, src=(UPat(Ops.UNSHARD, name="dest"), UPat(Ops.UNSHARD, name="src"))))), store_after_multi), (UPat(Ops.COPY, src=(UPat(Ops.UNSHARD, name="multi"),), name="copy"), lambda multi,copy: copy_multi(multi, copy.arg)), (UPat(Ops.ALLREDUCE, src=(UPat(Ops.UNSHARD, name="multi"),), name="red"), - lambda multi,red: multi.src[0].allreduce(*red.arg).unshard(multi.axis, multi.src[1])), + lambda multi,red: multi.src[0].allreduce(*red.arg).unshard(multi.arg, multi.src[1:])), # resolve TUPLE+GETTUPLE (needed in multi) (UPat(Ops.GETTUPLE, src=(UPat(Ops.TUPLE, name="t"),), name="g"), lambda g,t: t.src[g.arg]), # GETTUPLE on UNSHARD: passthrough UNSHARD (e.g. when FUNCTION was replaced by UNSHARD(GETTUPLE(...))) (UPat(Ops.GETTUPLE, src=(UPat(Ops.UNSHARD, name="multi"),), name="g"), - lambda g, multi: multi.src[0].gettuple(g.arg).unshard(multi.axis, multi.src[1]) if multi.src[0].op in {Ops.FUNCTION, Ops.TUPLE} - else multi), + lambda g, multi: multi.src[0].gettuple(g.arg).unshard(multi.arg, multi.src[1:]) if multi.src[0].op in {Ops.FUNCTION, Ops.TUPLE} else multi), # rewrite into FUNCTION calls explicitly for UNSHARD (value-producing) (UPat(Ops.FUNCTION, name="call"), rewrite_into_function), (UPat((Ops.CALL, Ops.FUNCTION, Ops.AFTER), src=(UPat(Ops.UNSHARD, name="multi"), ), name="root", allow_any_len=True), passthrough_multi), diff --git a/tinygrad/uop/ops.py b/tinygrad/uop/ops.py index 5c446cf989..a917297e06 100644 --- a/tinygrad/uop/ops.py +++ b/tinygrad/uop/ops.py @@ -439,7 +439,7 @@ class UOp(RandMixin, metaclass=UOpMetaClass): case Ops.FLIP: if len(ps) != len(self.marg) or not all(isinstance(x, bool) for x in self.marg): raise ValueError(f"bad flip on {ps}, {self.marg}") return ps - case Ops.UNSHARD: return tuple(s*(int(self.src[1].vmax)+1) if a == self.axis else s for a,s in enumerate(ps)) + case Ops.UNSHARD: return tuple(s*(int(self.src[1:][self.arg.index(a)].vmax)+1) if a in self.arg else s for a,s in enumerate(ps)) case Ops.REDUCE: num_axes = self.arg[1] if not isinstance(num_axes, int) or num_axes < 0 or num_axes > len(ps): @@ -665,15 +665,24 @@ class UOp(RandMixin, metaclass=UOpMetaClass): # *** multi-device helpers *** - def unshard(self, axis:int|None, device_range:UOp|None=None): + def unshard(self, axis:int|tuple[int, ...]|None, device_range:UOp|tuple[UOp, ...]|None=None): assert axis is not None, "multi None is no longer supported" - # an UNSHARD always has two srcs: the value and the range it ends (defaults to a DEVICE range over the devices) - # the range need not be DEVICE, e.g. a LOCAL range shards a kernel tile into per-thread fragments + # an UNSHARD carries the value and one sharding range per sharded axis (arg is the tuple of sharded axes, + # sorted). the single-axis axis form defaults the range to a DEVICE range over the devices; a range need not + # be DEVICE, e.g. a LOCAL range shards a kernel tile into per-thread fragments + if isinstance(axis, int): axis = (axis,) if device_range is None: assert isinstance(self.device, tuple), f"multi device must be tuple, {self.device} isn't" - device_range = UOp.range(len(self.device), -1, AxisType.DEVICE) - assert device_range.op is Ops.RANGE - return UOp(Ops.UNSHARD, src=(self, device_range), arg=axis) + device_range = (UOp.range(len(self.device), -1, AxisType.DEVICE),) + if isinstance(device_range, UOp): device_range = (device_range,) + assert isinstance(device_range, tuple) and len(axis) == len(device_range) and len(set(axis)) == len(axis) + axis, device_range = map(tuple, zip(*sorted(zip(axis, device_range)))) + return UOp(Ops.UNSHARD, src=(self, *device_range), arg=axis) + + @property + def sharding(self) -> tuple[tuple[int, UOp], ...]: + """(axis, RANGE) pairs this value is sharded over (the source of truth for shard bounds/counts).""" + return tuple(zip(self.arg, self.src[1:])) if self.op is Ops.UNSHARD else () @property def bounds(self): @@ -685,7 +694,9 @@ class UOp(RandMixin, metaclass=UOpMetaClass): def axis(self) -> int|None: # COPY removes axis. TODO: add more tests for this, and consider MSELECT/MSTACK if self.op is Ops.COPY: return None - if self.op is Ops.UNSHARD: return self.arg + if self.op is Ops.UNSHARD: + if len(self.arg) != 1: raise RuntimeError(f"UOp is sharded on multiple axes {self.arg}, use .sharding") + return self.arg[0] # GETTUPLE: axis comes from the specific TUPLE element, not src[0] if self.op is Ops.GETTUPLE: in_tuple = self.src[0].src[0] if self.src[0].op is Ops.FUNCTION else self.src[0] @@ -723,7 +734,6 @@ class UOp(RandMixin, metaclass=UOpMetaClass): return self.pad(tuple((0,0) if a != axis else (bsz*dnum, bsz*(dcount-1) - bsz*dnum) for a in range(len(self.shape)))) def _shard(self, axis:int, rng:UOp) -> UOp: - assert rng.op is Ops.RANGE, f"_shard requires a RANGE, got {rng.op}" if len(self.shape) == 0: return self # scalars broadcast, no sharding needed dcount = int(rng.vmax)+1 if self.shape[axis] % dcount != 0: raise RuntimeError(f"multi axis uneven: {self.shape[axis]=} {axis=} {dcount=}") @@ -879,7 +889,7 @@ class UOp(RandMixin, metaclass=UOpMetaClass): # CL 1.1 provides the clCreateSubBuffer API, but at the time of writing, relevant CL runtimes (rusticl, adreno, nvidia, amd) do not provide # reasonable values for CL_DEVICE_MEM_BASE_ADDR_ALIGN. cl_ext_buffer_device_address could potentially help, but this extension is not provided # by relevant CL runtimes at time of writing. - if any(d.startswith(("WEBGPU", "CL")) for d in ((self.device,) if isinstance(self.device, str) else self.device)): return None + if (dev:=self.device) is not None and any(d.startswith(("WEBGPU", "CL")) for d in ((dev,) if isinstance(dev, str) else dev)): return None idx = self.flatten().index(UOp.range(self.numel(), 0)) out = graph_rewrite(idx, pm_mops+symbolic+pm_contiguous_view_offset, ctx=self, name="contiguous_view_offset") diff --git a/tinygrad/uop/spec.py b/tinygrad/uop/spec.py index 1596526d9b..5298129d0a 100644 --- a/tinygrad/uop/spec.py +++ b/tinygrad/uop/spec.py @@ -176,9 +176,9 @@ spec_tensor = PatternMatcher([ len(red.arg) == 2 and red.arg[0] in GroupOp.Reduce and is_device(red.arg[1])), # UNSHARD/MSELECT/MSTACK - # an UNSHARD always has two srcs: the value and the range it ends (usually DEVICE, but any typed range can be sharded over) - (UPat(Ops.UNSHARD, name="multi"), lambda multi: len(multi.src) == 2 and matches_dtype(multi.src[0], multi.dtype) - and isinstance(multi.arg, int) and multi.src[1].op is Ops.RANGE and isinstance(multi.src[1].arg[-1], AxisType)), + # an UNSHARD carries the value and one sharding range per sharded axis (usually a DEVICE RANGE, but can be a derived expression) + (UPat(Ops.UNSHARD, name="multi"), lambda multi: len(multi.src) == 1+len(multi.arg) and matches_dtype(multi.src[0], multi.dtype) + and all(isinstance(a, int) for a in multi.arg) and all(r.dtype in dtypes.weaks for r in multi.src[1:])), (UPat(Ops.MSELECT, name="x"), lambda x: isinstance(x.src[0].device, tuple) and x.arg < len(x.src[0].device)), (UPat(Ops.MSTACK, name="x"), lambda x: all(isinstance(s.device, str) for s in x.src) or (all_same(x.src) and x.src[0].device is None)),