Mj|_hddlmZmZddlmZddlmZmZmZddl m Z ddl Z GddZ ed Gd d e Z ed Gd d e Zed Gdde Zed Gdde Zed Gdde Zed Gdde ZeddGdde ZGddZe dZed GddeZeddGddeZeddGd d!eZed Gd"d#eZd&d$Zd%ZdS)') dataclassfield)List)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functionNc>eZdZdZedZedZdS)DistributedLayoutz@ Base class for distributed memory layouts in Gluon IR. c t|SNrselfs p/home/longshao/multi-rider-rag/.venv/lib/python3.11/site-packages/triton/experimental/gluon/language/_layouts.pytypezDistributedLayout.typed###c td)Nz-DistributedLayout subclasses must define rank)NotImplementedErrorrs rrankzDistributedLayout.ranks!"QRRRrN)__name__ __module__ __qualname____doc__propertyrrrrr r sY$$X$SSXSSSrr T)frozenc0eZdZdZdZedZdS) AutoLayoutc*|Sr )get_auto_layoutrbuilders r_to_irzAutoLayout._to_irs&&(((rcdS)NALrrs rmanglezAutoLayout.mangletrc td)NzAutoLayout has no rank ValueErrorrs rrzAutoLayout.rank s1222rNrrrr%r(rrrrrr r sM)))33X333rr c0eZdZdZdZedZdS)CoalescedLayoutc*|Sr )get_coalesced_layoutr#s rr%zCoalescedLayout._to_ir(s++---rcdS)NCLrrs rr(zCoalescedLayout.mangle+r)rc td)NzCoalescedLayout has no rankr+rs rrzCoalescedLayout.rank.s6777rNr-rrrr/r/%sM...88X888rr/ceZdZUdZeeed<eeed<eeed<eeed<eeZ eeeed<fdZ d Z d e fd Z d Zed ZxZS) BlockedLayouta Represents a blocked layout, partitioning a tensor across threads, warps, and CTAs. Args: size_per_thread (List[int]): Number of elements per thread per dimension. threads_per_warp (List[int]): Number of threads per warp per dimension. warps_per_cta (List[int]): Number of warps per CTA per dimension. order (List[int]): The ordering of dimensions for partitioning. cga_layout (Optional[List[List[int]]]): Bases describing how CTAs tile each dimension. size_per_threadthreads_per_warp warps_per_ctaorderdefault_factory cga_layoutctdt|jtdt|jtdt|jtdt|jt|j}t|d|j t|j|ksJt|j|ksJt|j|ksJdS)Nr7r8r9r:r=) super __setattr__rr7r8r9r:lenobjectr=)rr __class__s r __post_init__zBlockedLayout.__post_init__Es -/CDDX/Y/YZZZ .0DTEZ0[0[\\\ O-A$BT-U-UVVV G%9$*%E%EFFF4'((4t???4())T11114%&&$....4:$&&&&&&rcf||j|j|j|j|jSr )get_blocked_layoutr7r8r9r:r=r#s rr%zBlockedLayout._to_irQs6))   !   J O    rreturnc d}||j}||j}||j}||j}|jr$dd|jDnd}d|d|d|d|d|d S)NcZ|dSdtt|SN_joinmapstrxs r stringifyz'BlockedLayout.mangle..stringify\&yr88CQKK(( (rrLc3fK|],}dtt|V-dS~NrM.0vecs r z'BlockedLayout.mangle..e6QQ#chhs3}}55QQQQQQrrKB)r7r8r9r:r=rN)rrSr7r8r9r:r=s rr(zBlockedLayout.mangleZs ) ) ) $)D$899$9T%:;;! $"455  $*%%UYUdlSXXQQQQQQQQjl ]?]]%5]] ]]]]PZ]]]]rc tt|jt|jt|jt|jtd|jDfS)Nc34K|]}t|VdSr tuplerXs rr[z)BlockedLayout.__hash__..js(-T-TSeCjj-T-T-T-T-T-Tr)hashrar7r8r9r:r=rs r__hash__zBlockedLayout.__hash__hsnU4/00%8M2N2NPUVZVhPiPi4:&&-T-TDO-T-T-T(T(TVWW Wrc*t|jSr )rAr:rs rrzBlockedLayout.rankl4:rrrrrrint__annotations__rlistr=rDr%rPr(rcrr __classcell__rCs@rr6r63s  #Y3i9 9"'%"="="=JT#Y=== ' ' ' ' '    ^ ^ ^ ^ ^WWWXrr6c|eZdZUdZeed<eed<fdZdZde fdZ dZ e d Z e d ZxZS) SliceLayoutz Represents a layout corresponding to slicing a distributed tensor along one dimension. Args: dim (int): The dimension index to slice. parent (DistributedLayout): The parent layout before slicing. dimparentctdt|jtdt|jdS)Nrnro)r?r@rrnrorrCs rrDzSliceLayout.__post_init__}sS E#7#A#ABBB H&:4;&G&GHHHHHrch||j|j|Sr )get_slice_layoutrnror%r#s rr%zSliceLayout._to_irs2'' H K  w ' '   rrGcLd|jd|jdS)NSLrL)rnror(rs rr(zSliceLayout.mangles+7DH77t{11337777rc8t|j|jfSr )rbrnrors rrczSliceLayout.__hash__sTXt{+,,,rc |jjdz SNrorrs rrzSliceLayout.ranks{!##rcjj}|sgSjj}djcxkr|ksnJfd|DS)NrcTg|]$}|dj|jdzdz%Srx)rn)rYbasisrs r z*SliceLayout.cga_layout..s8WWWEitxi 5A#77WWWr)ror=rrn)rparent_cga_layoutrs` rr=zSliceLayout.cga_layoutsg K2  I{DH####t######WWWWEVWWWWrrrrrrgrhr rDr%rPr(rcrrr=rjrks@rrmrmqs HHH IIIII   88888---$$X$XXXXXXXXrrmceZdZUdZeeeed<eeeed<eeeed<eeeed<eeed<fdZdZd Z d Z e d Z xZ S) DistributedLinearLayouta Represents a linear distributed layout with explicit bases at register, lane, warp, and block levels. See: https://arxiv.org/abs/2505.23819 for reference. Args: reg_bases (List[List[int]]): Bases for register-level distribution. lane_bases (List[List[int]]): Bases for lane-level distribution. warp_bases (List[List[int]]): Bases for warp-level distribution. block_bases (List[List[int]]): Bases for block-level distribution. shape (List[int]): The tensor global shape. reg_bases lane_bases warp_bases block_basesshapec0tdt|jtdt|jtdt|jtdt|jtdt|jt|j}|jD]}t||ksJ|jD]}t||ksJ|jD]}t||ksJ|jD]}t||ksJdS)Nrrrrr) r?r@rrrrrrrArrr}rCs rrDz%DistributedLinearLayout.__post_init__s| Kt~)F)FGGG L-*H*HIII L-*H*HIII M=9I+J+JKKK G]4:%>%>???4:^ & &Eu::%%%%%_ & &Eu::%%%%%_ & &Eu::%%%%%% & &Eu::%%%%% & &rcf||j|j|j|j|jSr )get_distributed_linear_layoutrrrrrr#s rr%zDistributedLinearLayout._to_irs644T^T_VZVegkgw59ZAA Arc Xd|jd|jd|jd|jd|jd S)NDLLrL)rrrrrrs rr(zDistributedLinearLayout.mangles@lT^lldollll$JZll]a]gllllrc vtttt|jttt|jttt|jttt|jt|jfSr )rbrarOrrrrrrs rrcz DistributedLinearLayout.__hash__s #eT^,, - - #eT_-- . . #eT_-- . . #eT-.. / / $*       rc*t|jSr )rArrs rrzDistributedLinearLayout.rankrer)rrrrrrgrhrDr%r(rcrrrjrks@rrrs  DIT#YT#Yd3i   9&&&&&$AAAmmm   XrrceZdZUdZeed<eed<eed<fdZdZde fdZ d Z e d Z e d ZxZS) DotOperandLayouta  Represents a layout for a dot operand. Args: operand_index (int): 0 for LHS and 1 for RHS of the dot operation. parent (DistributedLayout): The parent layout, representing the MMA. k_width (int): Number of elements per 32-bits. operand_indexrok_widthc@tdt|jtdt|jtdt|jdS)Nrror)r?r@rrrorrqs rrDzDotOperandLayout.__post_init__sy O-A$BT-U-UVVV H&:4;&G&GHHH I';DL'I'IJJJJJrct||j|j||jSr )get_dot_operand_layoutrror%rr#s rr%zDotOperandLayout._to_irs2--d.@$+BTBTU\B]B]_c_klllrrGc\d|jd|jd|jdS)NDOrL)rror(rrs rr(zDotOperandLayout.mangles6PD&PP););)=)=PP PPPPrcDt|j|j|jfSr )rbrrorrs rrczDotOperandLayout.__hash__sT'dlCDDDrc|jjSr rzrs rrzDotOperandLayout.ranks {rcXtt|jdgpg}|sgS|jjt fd|DsJ|jdkrdz ndz }d|cxkrksnJg}|D]+}t |}d||<||,|S)Nr=c3>K|]}t|kVdSr rA)rYr}rs rr[z.DotOperandLayout.cga_layout..s.EE%3u::%EEEEEErrry)rgetattrrorallrriappend)rrk_dimderivedr} new_basisrs @rr=zDotOperandLayout.cga_layouts0lTV1W1WXX^\^  I{EEEE3DEEEEEEEE .!33qE    D      & & &EU I Ie  NN9 % % % %rrrks@rrrs  LLLKKKKK mmmQQQQQEEE  X Xrr)reqceZdZUdZeeed<eeed<eeed<eeZ eeeed<fdZ dZ d e fd Z d Zed ZxZS) NVMMADistributedLayouta_ Represents a layout for NVIDIA MMA (tensor core) operations. Args: version (List[int]): Version identifier for the MMA instruction. warps_per_cta (List[int]): Number of warps per CTA. instr_shape (List[int]): Instruction shape for MMA. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. versionr9 instr_shaper;r=ctdt|jtdt|jtdt|jt |d|jdS)Nrr9rr=)r?r@rrr9rrBr=rqs rrDz$NVMMADistributedLayout.__post_init__s I';DL'I'IJJJ O-A$BT-U-UVVV M+?@P+Q+QRRR4t?????rcZ||j|j|j|jSr )get_mma_layoutrr9r=rr#s rr%zNVMMADistributedLayout._to_ir!s0%% L   O      rrGc |jr$dd|jDnd}d|jd|jd|jd|d S)NrLc3fK|],}dtt|V-dSrVrMrXs rr[z0NVMMADistributedLayout.mangle..*r\rrKMMA__MMA)r=rNrr9rrr=s rr(zNVMMADistributedLayout.mangle)scUYUdlSXXQQQQQQQQjl ]dl]]T%7]]$:J]]Z]]]]rc tt|jt|jt|jtd|jDfS)Nc34K|]}t|VdSr r`rXs rr[z2NVMMADistributedLayout.__hash__../(AA#5::AAAAAAr)rbrarr9rr=rs rrczNVMMADistributedLayout.__hash__-s_U4<((%0B*C*CU4K[E\E\AAAAAAACDD Drc*t|jSr )rAr9rs rrzNVMMADistributedLayout.rank1s4%&&&rrfrks@rrr s#Y9c"'%"="="=JT#Y===@@@@@   ^^^^^DDD''X'''''rrc(eZdZdZedZdS) SharedLayoutz; Base class for shared memory layouts in Gluon IR. c t|Sr rrs rrzSharedLayout.type;rrN)rrrrrrrrrrr6s9$$X$$$rrc|s|St|}t|d}dg|z}|D]H}t||ksJt|D]!}t||||||<"It|D]}||xxdzcc<t|D]=}||||zdksJd|d|||xx||zcc<>|S)NrryrzShape z is not divisible by CGA layout )rirArangemax)rr= shape_per_ctar cga_shaper}irns r_get_shape_per_ctar@s9  KKM z!}  Dd I775zzT!!!!t 7 7Ay|U1X66IaLL 74[[!  T{{..S!IcN2a7779u%9u9uis9u9u777cy~- rceZdZUdZeed<eed<dZeed<dZeed<dZ eed<e e Z e e eed <fd Zd ZeeddZdefdZdZxZS)NVMMASharedLayouta Represents a layout for shared memory suitable for NVIDIA MMA operations. Args: swizzle_byte_width (int): Width in bytes for swizzling. element_bitwidth (int): Bitwidth of element type. rank (int): Rank of the tensor. transposed (bool): Whether the layout is transposed. fp4_padded (bool): Whether FP4 padding is used. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. swizzle_byte_widthelement_bitwidthrrF transposed fp4_paddedr;r=ctdt|jtdt|jtdt|jtdt|j|jpg}|r t|d|j ksJtdt|j tdt||jdvsJ|jd vsJdS) Nrrrrrrr=) @)rrr) r?r@rrrrrr=rAr)rr=rCs rrDzNVMMASharedLayout.__post_init__hsG 02FtG^2_2_``` .0DTEZ0[0[\\\ L*>t*O*OPPP L*>t*O*OPPP_*  3z!}%%2222 F$8$C$CDDD L*>z*J*JKKK$7777&*:::::::rcr||j|j|j|j|j|jSr )get_nvmma_shared_layoutrrrrr=rr#s rr%zNVMMASharedLayout._to_irys:..  #  ! O O O I    rNc|rdnd}||nt||}t|}|r|dd|ddz}|d|z}||jzdz} | dkr | dzdkrd} n&| dkr | dzdkrd} n| d kr | d zdkrd } nd} d} |ddD]} | | z} t|dks| dkrd} t| |j|||| S) zReturns an NVMMASharedLayout with default swizzling for a given shape. This picks the largest swizzle pattern compatible with the shape, which allows emitting the fewest TMA or MMA messages. rryNrrrrr)rrrrrr=)rrAprimitive_bitwidthr) block_shapedtyperrr=packing_factorrrcontig_dim_sizecontig_dim_bytesrflatten_outer_dimsizes rget_default_forz!NVMMASharedLayout.get_default_forso)/a'1'9 ?QR]_i?j?j ;  B)!""- bqb0AAM'+n<*U-EEJ s " "'7#'='B'B!$    # #(82(=(B(B!#    # #(82(=(B(B!#  !" !#2#& & &D  %   {  a  #4q#8#8!"  1"5!!!     rrGc |jr$dd|jDnd}d|jd|jd|jd|jd|d S)NrLc3fK|],}dtt|V-dSrVrMrXs rr[z+NVMMASharedLayout.mangle..r\rrKNVMMA__NVMMA)r=rNrrrrrs rr(zNVMMASharedLayout.manglesUYUdlSXXQQQQQQQQjl A/AA$2GAA$/AA\`\kAAnxAAA Arc t|j|j|j|j|j|jrtd|jDndfS)Nc34K|]}t|VdSr r`rXs rr[z-NVMMASharedLayout.__hash__..rr)rbrrrrrr=rars rrczNVMMASharedLayout.__hash__s`T,d.CTYPTP_aeapEI_^UAAAAAAAAZ^`aa ar)FFN)rrrrrgrhrrboolrrrir=rrDr% staticmethodr rrPr(rcrjrks@rrrTs!  D#MMMJJ"'%"="="=JT#Y===;;;;;"   # # # \# JAAAAAaaaaaaarrceZdZUdZeed<eed<eed<eeed<eeZ eeeed<fdZ d Z d e fd Z d ZxZS) SwizzledSharedLayoutaq Represents a generic swizzled shared memory layout. Args: vec (int): Vector width for swizzling. per_phase (int): Elements per swizzle phase. max_phase (int): Maximum number of swizzle phases. order (List[int]): Dimension ordering for swizzling. cga_layout (Optional[List[List[int]]]): Bases describing CTA tiling. rZ per_phase max_phaser:r;r=ctdt|jtdt|jtdt|jtdt|jt|d|jdS)NrZrrr:r=) r?r@rrZrrr:rBr=rqs rrDz"SwizzledSharedLayout.__post_init__s E#7#A#ABBB K)=dn)M)MNNN K)=dn)M)MNNN G%9$*%E%EFFF4t?????rcf||j|j|j|j|jSr )get_swizzled_shared_layoutrZrrr:r=r#s rr%zSwizzledSharedLayout._to_irs311 H N N J O    rrGc d}|jr$dd|jDnd}d|jd|jd|jd||jd|d S)NcZ|dSdtt|SrJrMrQs rrSz.SwizzledSharedLayout.mangle..stringifyrTrrLc3fK|],}dtt|V-dSrVrMrXs rr[z.SwizzledSharedLayout.mangle..r\rrKSSS__SSS)r=rNrZrrr:)rrSr=s rr(zSwizzledSharedLayout.mangles ) ) ) VZUdlSXXQQQQQQQQjl kdhkkkk$.kk99TZCXCXkk[ekkkkrc t|j|j|jt |jt d|jDfS)Nc34K|]}t|VdSr r`rXs rr[z0SwizzledSharedLayout.__hash__..s-OvOv_bPUVYPZPZOvOvOvOvOvOvr)rbrZrrrar:r=rs rrczSwizzledSharedLayout.__hash__sP Xt~t~uTZ7H7H%OvOvfjfuOvOvOvJvJv wyy yr)rrrrrgrhrrrir=rDr%rPr(rcrjrks@rrrs   HHHNNNNNN 9"'%"="="=JT#Y===@@@@@   lllllyyyyyyyrrceZdZUdZeeeed<eeeed<eeeed<eeed<fdZdZde fd Z d Z e e d Zd ZxZS) PaddedSharedLayouta Represents a layout for the access to shared memory. Compared to SwizzledSharedLayout, it combined padding and element reordering via linear transformation (e.g. row permutation) to avoid shared memory bank conflicts. After every interval tensor elements, the corresponding number of padding elements are inserted. If a position corresponds to multiple intervals, the padding amounts are summed. In the following example of a tensor, `eM` represents original elements in the and `pN` represents padded element. Before padding, the shared memory looks like: [e0, e1, e2, e3, e4, e5, e6, e7, ...] After padding with interval-padding list [[2, 1], [4, 2]] with an identity remapping, the shared memory will be [e0, e1, p0, e2, e3, p1, p2, p3, e4, e5, p4, e6, e7, p5, p6, p7, ...] Furthermore this encoding allows for a linear remapping from the 1-D shared memory offset to logical n-D tensor elements. The remapping is given in the form of linear bases mapping from offset to [dim0, dim1...dimN-1]. See LinearLayout.h for more details how linear layouts are applied to remap elements. Some concrete examples using `xN` and `yN` to mean the logical n-D tensor elements and `pN` to mean padding: After padding for shape = [8] with interval-padding list [[2, 2]], offset_bases = [[2], [1]] and block_bases = []: [x0, x2, p0 p1, x1, x3] After padding for shape = [8, 4] with interval_padding_pairs = [[8, 1]], offset_bases = [[0, 1], [0, 2], /*gap, stride by 2 rows*/[2, 0], [4, 0], [1, 0]]] and block_bases = []: [ x0y0, x0y1, x0y2, x0y3, x2y0, x2y1, x2y2, x2y3, p0, x4y0, x4y1, x4y2, x4y3, x6y0, x6y1, x6y2, x6y3, p1, x1y0, x1y1, x1y2, x1y3, x3y0, x3y1, x3y2, x3y3, p2, x5y0, x5y1, x5y2, x5y3, x7y0, x7y1, x7y2, x7y3, ] Args: interval_padding_pairs (List[int]): List of [interval, padding] pair and both interval and padding must be powers of 2. offset_bases (List[int]): Bases for shared memory offsets block_bases (List[List[int]]): Bases for block-level shared memory offsets. shape (List[int]): n-D logical shared memory shape interval_padding_pairs offset_basesrrcttdt|jtdt|jtdt|jtdt|jt|j}|jD]}t||ksJ|jD]}t||ksJ|dS)Nrrrr) r?r@rrrrrrAverifyrs rrDz PaddedSharedLayout.__post_init__%s  4mDD_6`6`aaa NM$:K,L,LMMM M=9I+J+JKKK G]4:%>%>???4:& & &Eu::%%%%%% & &Eu::%%%%% rctt|j\}}||||j|j|jSr )ziprget_padded_shared_layoutrrr)rr$ intervalspaddingss rr%zPaddedSharedLayout._to_ir4s;!4#>? 8// 8TEVX\XhjnjtuuurrGc Hd|jd|jd|jd|jd S)N PaddedShared_rL _PaddedShared)rrrrrs rr(zPaddedSharedLayout.mangle8s9}t:}}T=N}}QUQa}}dhdn}}}}rc|j}t|dks Jdtd|DsJt|\}}t t |}t|t|ksJdtfd|Ds Jdtfd|Ds Jdt|j}|dks Jd dS) NrzVPaddedSharedLayout interval_padding_pairs must have at least one interval-padding pairc3<K|]}t|dkVdS)rNr)rYpairs rr[z,PaddedSharedLayout.verify..>s,44d3t99>444444rc&|dko ||dz zdkSNrryrns rz+PaddedSharedLayout.verify..D!a%".Es-77==##777777rz;PaddedSharedLayout interval values must all be power of twoc3.K|]}|VdSr rrs rr[z,PaddedSharedLayout.verify..Fs-66==##666666rz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty)rrArrrisetr)rpairsrrunique_intervalsrrs @rrzPaddedSharedLayout.verify;s+5zzA~~~w~~~44e44444444!5k 8I//#$$I6666<< 7777Y77777vv9vvv76666X66666tt8ttt64:axxxExxxxxrc t|t|ksJdtfd|DsJt|}g}|D]httt j|D]1|fdt|D2it||g|S)zReturns a PaddedSharedLayout with the given interval and padding pairs and an identity mapping as the linear component for the given shape and order. c&|dko ||dz zdkSrrrs rrz6PaddedSharedLayout.with_identity_for..Qrrc3.K|]}|VdSr rrs rr[z7PaddedSharedLayout.with_identity_for..Rs-33==##333333rc*g|]}|krdzndS)ryrr)rYrr}rns rr~z8PaddedSharedLayout.with_identity_for..Ys)$X$X$Xq188Q%ZZ$X$X$Xr)rArrrgmathlog2rr)rrr:rrr}rnrs @@@rwith_identity_forz$PaddedSharedLayout.with_identity_forKs 5zzSZZ''''<< 3333U333333335zz  Z ZCs49U3Z#8#899:: Z Z##$X$X$X$X$XERVKK$X$X$XYYYY Z""8,ERRRrc *tttt|jttt|jttt|jt|jfSr )rbrarOrrrrrs rrczPaddedSharedLayout.__hash__]spU3ud&ABBCCU3uVZVgKhKhEiEi3ud&67788% :K:KMNN Nr)rrrrrrgrhrDr%rPr(rrr rrcrjrks@rrrs88r!cO+++tCy/!!!d3i   9     vvv~~~~~FFF SS\S NNNNNNNrrceZdZUdZeeeed<eeZ eeeed<dZ eed<fdZ dZ d e fd Zd ZxZS) SharedLinearLayoutzGRepresents a shared memory layout defined via an explicit LinearLayout.rr;rr alignmentctdt|jtdt|jtdt |jt|jdks Jdt|jd}|dks Jd|jD]}t||ksJ|jD]}t||ksJ|jdkr|j|jdz zdks JddS)Nrrrrz1SharedLinearLayout offset_bases must not be emptyryz#>#>Kd3i>>>Is K K K K Keeeddddd       rrcdg|z}|s|Sd}|D]T}tdt|Dd}||}||xxdzcc<>|s|J||xxdzcc<U|S)Nryc3,K|]\}}|dk |VdS)rNr)rYrvs rr[z bases_per_dim..s*==$!Qa1ffAffff==rr)next enumerate)basesrskip_broadcastresult non_zero_idxr}idxs r bases_per_dimr)sS4ZF  L & &==)E"2"2===tDD ?L 3KKK1 KKKK &+++ < A % Mrct|tr"t|jt |St|t t frt|j|S|jSr ) isinstancerr)rrArmrr9ro)layoutrs rr9r9sa&122$V.E ;;; F[*:; < <$V]E222##r)T) dataclassesrrtypingrtriton.language.corerrrtriton.runtime.jitr rr r r/r6rmrrrrrrrrrr)r9rrrr1s;((((((((TTTTTTTTTT111111  S S S S S S S S $ 3 3 3 3 3" 3 3 3 $ 8 8 8 8 8' 8 8 8 $:::::%:::z $'X'X'X'X'X#'X'X'XT $66666/666r $00000(000f $4   ('('('('('.('('! ('V$$$$$$$$& $[a[a[a[a[a [a[a[a| $4   .y.y.y.y.y<.y.y! .yb $4   yNyNyNyNyNyNyN! yNx $! ! ! ! ! ! ! ! J,$$$$$r