~i|_NSSKJrJr SSKJr SSKJrJrJr SSK J r SSK r "SS5r \"SS 9"S S \ 55r \"SS 9"S S \ 55r\"SS 9"SS\ 55r\"SS 9"SS\ 55r\"SS 9"SS\ 55r\"SS 9"SS\ 55r\"SSS9"SS\ 55r"SS5r\ S5r\"SS 9"SS\55r\"SSS9"SS\55r\"SSS9"S S!\55r\"SS 9"S"S#\55rS&S$jrS%rg)') dataclassfield)List)_unwrap_if_constexpr _unwrap_shapeconstexpr_type)constexpr_functionNc8\rSrSrSr\S5r\S5rSrg)DistributedLayout z8 Base class for distributed memory layouts in Gluon IR. c[U5$Nrselfs h/home/mande/repo/quber/.venv/lib/python3.13/site-packages/triton/experimental/gluon/language/_layouts.pytypeDistributedLayout.type d##c[S5e)Nz-DistributedLayout subclasses must define rank)NotImplementedErrorrs rrankDistributedLayout.ranks!"QRRrN) __name__ __module__ __qualname____firstlineno____doc__propertyrr__static_attributes__rrrr r s1$$SSrr T)frozenc0\rSrSrSrSr\S5rSrg) AutoLayoutc"UR5$r)get_auto_layoutrbuilders r_to_irAutoLayout._to_irs&&((rcg)NALrrs rmangleAutoLayout.manglerc[S5e)NzAutoLayout has no rank ValueErrorrs rrAutoLayout.rank s122rrN rrrrr+r/r!rr"rrrr%r%s )33rr%c0\rSrSrSrSr\S5rSrg)CoalescedLayout%c"UR5$r)get_coalesced_layoutr)s rr+CoalescedLayout._to_ir(s++--rcg)NCLrrs rr/CoalescedLayout.mangle+r1rc[S5e)NzCoalescedLayout has no rankr3rs rrCoalescedLayout.rank.s677rrNr6rrrr8r8%s .88rr8c^\rSrSr%Sr\\\S'\\\S'\\\S'\\\S'\"\ S9r \\\\S'U4S jr S r S \ 4S jrS r\S5rSrU=r$) BlockedLayout3a 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_layoutc,>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [UR5n[RUSUR5 [UR5U:Xde[UR 5U:Xde[UR 5U:Xdeg)NrErFrGrHrK) super __setattr__rrErFrGrHlenobjectrK)rr __class__s r __post_init__BlockedLayout.__post_init__Es -/CDDXDX/YZ .0DTEZEZ0[\ O-A$BTBT-UV G%9$**%EF4''(4t?4(()T1114%%&$...4::$&&&rcURURURURURUR 5$r)get_blocked_layoutrErFrGrHrKr)s rr+BlockedLayout._to_irQs>))   ! !    JJ OO   rreturnc &SnU"UR5nU"UR5nU"UR5nU"UR5nUR(a"SR SUR55OSnSUSUSUSUSUS3 $)NcHUcgSR[[U55$N_joinmapstrxs r stringify'BlockedLayout.mangle..stringify\y88CQK( (rr\c3`# UH$nSR[[U55v M& g7f~Nr].0vecs r 'BlockedLayout.mangle..e"Q#chhs3}55,.r[B)rErFrGrHrKr^)rrcrErFrGrHrKs rr/BlockedLayout.mangleZs ) $D$8$89$T%:%:;!$"4"45 $**%UYUdUdSXXQQQjl ?#1%5$6a awaPZ|[\]]rc [[UR5[UR5[UR5[UR 5[SUR 5545$)Nc38# UHn[U5v M g7frtupleris rrl)BlockedLayout.__hash__..js-TOSeCjjO)hashrurErFrGrHrKrs r__hash__BlockedLayout.__hash__hs]U4//0%8M8M2NPUVZVhVhPi4::&-TDOO-T(TVW Wrc,[UR5$r)rOrHrs rrBlockedLayout.rankl4::rrrrrrr rint__annotations__rlistrKrRr+r`r/ryr!rr" __classcell__rQs@rrCrC3s{ #Y3i9 9"'"=JT#Y= '  ^ ^WrrCc|^\rSrSr%Sr\\S'\\S'U4SjrSr S\ 4Sjr S r \ S 5r\ S 5rS rU=r$) SliceLayoutqz 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. dimparentc>[TU]S[UR55 [TU]S[UR55 g)Nrr)rMrNrrrrrQs rrRSliceLayout.__post_init__}s5 E#7#AB H&:4;;&GHrclURURURRU55$r)get_slice_layoutrrr+r)s rr+SliceLayout._to_irs.'' HH KK  w '  rrWcXSURSURR5S3$)NSLr\)rrr/rs rr/SliceLayout.mangles)DHH:Qt{{1134B77rcD[URUR45$r)rxrrrs rrySliceLayout.__hash__sTXXt{{+,,rc4URRS- $)Nrrrs rrSliceLayout.ranks{{!##rcURRnU(d/$URRnSURs=::aU:de eUVs/sH#o3SURX0RS-S-PM% sn$s snfNrr)rrKrr)rparent_cga_layoutrbasiss rrKSliceLayout.cga_layouts{ KK22 I{{DHH#t#####EVWEVEitxx 5A#77EVWWWs*Brrrrrr rrr rRr+r`r/ryr!rrKr"rrs@rrrqs[ H I 88-$$XXrrc^\rSrSr%Sr\\\\S'\\\\S'\\\\S'\\\\S'\\\S'U4SjrS r S r S r \ S 5r S rU=r$)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_basesshapec>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [TU]S[UR55 [UR5nURHn[U5U:XaMe URHn[U5U:XaMe UR Hn[U5U:XaMe UR Hn[U5U:XaMe g)Nrrrrr) rMrNrrrrrrrOrrrrQs rrR%DistributedLinearLayout.__post_init__s  Kt~~)FG L-*HI L-*HI M=9I9I+JK G]4::%>?4::^^Eu:% %%$__Eu:% %%%__Eu:% %%%%%Eu:% %%&rcURURURURURUR 5$r)get_distributed_linear_layoutrrrrrr)s rr+DistributedLinearLayout._to_irs<44T^^T__VZVeVegkgwgw59ZZA Arc SURSURSURSURSURS3 $)NDLLr\)rrrrrrs rr/DistributedLinearLayout.manglesIT^^$Adoo%6a7H$JZJZI[[\]a]g]g\hhkllrc P[[[[UR55[[[UR55[[[UR 55[[[UR 55[UR545$r)rxrur_rrrrrrs rry DistributedLinearLayout.__hash__sn #eT^^, - #eT__- . #eT__- . #eT--. / $**     rc,[UR5$r)rOrrs rrDistributedLinearLayout.rankr}rr)rrrrr rrrrRr+r/ryr!rr"rrs@rrrsv DIT#YT#Yd3i 9&$Am rrc^\rSrSr%Sr\\S'\\S'\\S'U4SjrSr S\ 4S jr S r \ S 5r\ S 5rS rU=r$)DotOperandLayoutz 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_indexrk_widthc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 g)Nrrr)rMrNrrrrrs rrRDotOperandLayout.__post_init__sP O-A$BTBT-UV H&:4;;&GH I';DLL'IJrcURURURRU5UR5$r)get_dot_operand_layoutrrr+rr)s rr+DotOperandLayout._to_irs4--d.@.@$++BTBTU\B]_c_k_kllrrWcrSURSURR5SURS3$)NDOr\)rrr/rrs rr/DotOperandLayout.mangles6D&&'q););)=(>a ~RPPrcZ[URURUR45$r)rxrrrrs rryDotOperandLayout.__hash__s"T''dllCDDrc.URR$rrrs rrDotOperandLayout.ranks{{rc~^[[URS/55=(d /nU(d/$URRm[ U4SjU55(deUR S:XaTS- OTS- nSUs=::aT:de e/nUH#n[ U5nSXR'URU5 M% U$)NrKc3@># UHn[U5T:Hv M g7frrO)rjrrs rrl.DotOperandLayout.cga_layout..sE3D%3u:%3Dsrr)rgetattrrrallrrappend)rrk_dimderivedr new_basisrs @rrKDotOperandLayout.cga_layouts0lTV1WX^\^ I{{E3DEEEEE ..!3qE D     &EU I I  NN9 %'rrrrs@rrrsc  LK mQQE  rr)r#eqc^\rSrSr%Sr\\\S'\\\S'\\\S'\"\ S9r \\\\S'U4Sjr S r S \ 4S jrS r\S 5rSrU=r$)NVMMADistributedLayouti aC 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. versionrG instr_shaperIrKc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [ RUSUR5 g)NrrGrrK)rMrNrrrGrrPrKrs rrR$NVMMADistributedLayout.__post_init__sh I';DLL'IJ O-A$BTBT-UV M+?@P@P+QR4t?rczURURURURUR5$r)get_mma_layoutrrGrKrr)s rr+NVMMADistributedLayout._to_ir!s5%% LL    OO      rrWc UR(a"SRSUR55OSnSURSURSURSUS3 $)Nr\c3`# UH$nSR[[U55v M& g7frgr]ris rrl0NVMMADistributedLayout.mangle..*rnror[MMA__MMA)rKr^rrGrrrKs rr/NVMMADistributedLayout.mangle)sYUYUdUdSXXQQQjl dll^1T%7%7$8$:J:J9K1ZLX\]]rc [[UR5[UR5[UR5[SUR 5545$)Nc38# UHn[U5v M g7frrtris rrl2NVMMADistributedLayout.__hash__../A#5::rw)rxrurrGrrKrs rryNVMMADistributedLayout.__hash__-sNU4<<(%0B0B*CU4K[K[E\AAACD Drc,[UR5$r)rOrGrs rrNVMMADistributedLayout.rank1s4%%&&rrr~rs@rrr sr#Y9c"'"=JT#Y=@ ^^D''rrc(\rSrSrSr\S5rSrg) SharedLayouti6z3 Base class for shared memory layouts in Gluon IR. c[U5$rrrs rrSharedLayout.type;rrrN)rrrrr r!rr"rrrrr6s$$rrcU(dU$[U5n[US5nS/U-nUH8n[U5U:Xde[U5Hn[XFXV5XF'M M: [U5HnXF==S-ss'M [U5H+nX'XG-S:Xd SUSU35eX'==XG-ss'M- U$)NrrrzShape z is not divisible by CGA layout )rrOrangemax)rrK shape_per_ctar cga_shaperirs r_get_shape_per_ctar@s  KM z!} Dd I5zT!!!tAy|UX6IL 4[  T{!IN2a7u6%Hhisht9uu7y~- rc^\rSrSr%Sr\\S'\\S'Sr\\S'Sr\ \S'Sr \ \S '\ "\ S 9r \\\\S 'U4S jrS r\\SSj55rS\4SjrSrSrU=r$)NVMMASharedLayoutiTa 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_paddedrIrKcR>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 UR=(d /nU(a[US5UR:Xde[TU]S[UR55 [TU]S[U55 URS;deURS ;deg) NrrrrrrrK) @)rrr) rMrNrrrrrrKrOr)rrKrQs rrRNVMMASharedLayout.__post_init__hs 02FtG^G^2_` .0DTEZEZ0[\ L*>t*OP L*>t*OP__* z!}%2 22 F$8$CD L*>z*JK$$777&&*::::rcURURURURURUR UR 5$r)get_nvmma_shared_layoutrrrrrKrr)s rr+NVMMASharedLayout._to_irysC..  # #  ! ! OO OO OO II   rc U(aSOSnUcUO [X5n[U5nU(a USSUSS-nUSU-nXR-S-n U S:a U S-S:XaSn O&U S:a U S-S:XaSn OU S :a U S -S:XaS n OSn Sn USSHn X-n M [U5S:dU S:aSn [U URUUUUS 9$) 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. rrNrrrrr)rrrrrrK)rrOprimitive_bitwidthr) block_shapedtyperrrKpacking_factorrrcontig_dim_sizecontig_dim_bytesrflatten_outer_dimsizes rget_default_for!NVMMASharedLayout.get_default_fors)a'1'9 ?QR]?j ; )!"- bq0AAM'+n<*-E-EEJ s "'7#'='B!$   #(82(=(B!#   #(82(=(B!# !" !#2&D  % ' { a #4q#8!"  1"55!!!   rrWc UR(a"SRSUR55OSnSURSURSURSUR SUS3 $)Nr\c3`# UH$nSR[[U55v M& g7frgr]ris rrl+NVMMASharedLayout.mangle..rnror[NVMMA__NVMMA)rKr^rrrrrs rr/NVMMASharedLayout.manglesrUYUdUdSXXQQQjl //0$2G2G1H$//IZZ[\`\k\k[llmnxmyyA Arc [URURURURUR UR (a![SUR 5545$S45$)Nc38# UHn[U5v M g7frrtris rrl-NVMMASharedLayout.__hash__..rrw)rxrrrrrrKrurs rryNVMMASharedLayout.__hash__slT,,d.C.CTYYPTP_P_aeapapEI__UAAA`a aZ^`a arr)FFN)rrrrr rrrrboolrrrrKrrRr+ staticmethodr r r`r/ryr"rrs@rrrTs D#MJJ"'"=JT#Y=;" # # JAAaarrc^\rSrSr%Sr\\S'\\S'\\S'\\\S'\"\ S9r \\\\S'U4S jr S r S \ 4S jrS rSrU=r$)SwizzledSharedLayoutiaQ 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. rk per_phase max_phaserHrIrKc`>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [RUSUR5 g)NrkrrrHrK) rMrNrrkrrrHrPrKrs rrR"SwizzledSharedLayout.__post_init__s} E#7#AB K)=dnn)MN K)=dnn)MN G%9$**%EF4t?rcURURURURURUR 5$r)get_swizzled_shared_layoutrkrrrHrKr)s rr+SwizzledSharedLayout._to_irs811 HH NN NN JJ OO   rrWc SnUR(a"SRSUR55OSnSURSURSURSU"UR 5SUS3 $)NcHUcgSR[[U55$rZr]ras rrc.SwizzledSharedLayout.mangle..stringifyrerr\c3`# UH$nSR[[U55v M& g7frgr]ris rrl.SwizzledSharedLayout.mangle..rnror[SSS__SSS)rKr^rkrrrH)rrcrKs rr/SwizzledSharedLayout.manglesq ) VZUdUdSXXQQQjl dhhZq 0$..1A9TZZCXBYYZ[eZffjkkrc [URURUR[ UR 5[ SUR 5545$)Nc38# UHn[U5v M g7frrtris rrl0SwizzledSharedLayout.__hash__..sOvfu_bPUVYPZPZfurw)rxrkrrrurHrKrs rrySwizzledSharedLayout.__hash__sF XXt~~t~~uTZZ7H%OvfjfufuOvJv wy yrr)rrrrr rrrrrrKrRr+r`r/ryr"rrs@rrrs_  HNN 9"'"=JT#Y=@ llyyrrc^\rSrSr%Sr\\\\S'\\\\S'\\\\S'\\\S'U4SjrSr S \ 4S jr S r \ \S 55rS rSrU=r$)PaddedSharedLayoutia6 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_basesrrc>[TU]S[UR55 [TU]S[UR55 [TU]S[UR 55 [TU]S[UR 55 [UR 5nURHn[U5U:XaMe UR Hn[U5U:XaMe UR5 g)Nr2r3rr) rMrNrr2r3rrrOverifyrs rrR PaddedSharedLayout.__post_init__%s 4mDD_D_6`a NM$:K:K,LM M=9I9I+JK G]4::%>?4::&&Eu:% %%'%%Eu:% %%& rc[UR6up#URX#URURUR 5$r)zipr2get_padded_shared_layoutr3rr)rr* intervalspaddingss rr+PaddedSharedLayout._to_ir4s@!4#>#>? // TEVEVX\XhXhjnjtjtuurrWc pSURSURSURSURS3 $)N PaddedShared_r\ _PaddedShared)r2r3rrrs rr/PaddedSharedLayout.mangle8sCt::;1T=N=N# UHn[U5S:Hv M g7f)rNr)rjpairs rrl,PaddedSharedLayout.verify..>s4ed3t9>esc.US:=(a XS- -S:H$rrns r+PaddedSharedLayout.verify..D!a%"# UH nT"U5v M g7frrrjrG is_power_of_2s rrlrDEs7Y=##Yz;PaddedSharedLayout interval values must all be power of twoc34># UH nT"U5v M g7frrrLs rrlrDFs6X=##XrNz:PaddedSharedLayout padding values must all be power of twoz*PaddedSharedLayout order must not be empty)r2rOrr8rsetr)rpairsr:r;unique_intervalsrrMs @rr5PaddedSharedLayout.verify;s++5zA~www~4e44444!5k I/#$I666< 7Y777v9vv76X666t8tt64::axEEExrc ^[U5[U5:XdeSm[U4SjU55(de[U5n/nUHin[[[R "X555H8nUR [U5Vs/sHowU:XaSU-OSPM sn5 M: Mk [X/U5$s snf)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.US:=(a XS- -S:H$rrrFs rrH6PaddedSharedLayout.with_identity_for..QrJrc34># UH nT"U5v M g7frrrLs rrl7PaddedSharedLayout.with_identity_for..Rs3U=##UrNrr)rOrrrmathlog2rr1) r2rrHrr3rrrrMs @rwith_identity_for$PaddedSharedLayout.with_identity_forKs 5zSZ'''< 3U333335z Cs499UZ#89:##ERVK$XKq8Q%Z%BK$XY;""8ERR%YsC c  [[[[UR55[[[UR55[[[UR 55[UR 545$r)rxrur_r2r3rrrs rryPaddedSharedLayout.__hash__]s`U3ud&A&ABCU3uVZVgVgKhEi3ud&6&678% :KMN Nrr)rrrrr rrrrRr+r`r/r5rr r[ryr"rrs@rr1r1s8r!cO+tCy/!d3i 9 v~~F SS NNrr1c^\rSrSr%Sr\\\\S'\"\ S9r \\\\S'Sr \\S'U4Sjr S r S \4S jrS rS rU=r$)SharedLinearLayoutibzGRepresents a shared memory layout defined via an explicit LinearLayout.r3rIrr alignmentc`>[TU]S[UR55 [TU]S[UR55 [TU]S[ UR 55 [UR5S:wdS5e[URS5nUS:dS5eURHn[U5U:XaMe URHn[U5U:XaMe UR S:a UR UR S- -S:XdS5eg)Nr3rrarz1SharedLinearLayout offset_bases must not be emptyrzKd3i>Is Kedd  rr`cS/U-nU(dU$SnUHLn[S[U55S5nUbUnX6==S-ss'M2U(aM;UceX4==S-ss'MN U$)Nrc3:# UHupUS:wdM Uv M g7f)rNr)rjrvs rrl bases_per_dim..s="2$!a1fAA"2s  r)next enumerate)basesrskip_broadcastresult non_zero_idxridxs r bases_per_dimrxswS4ZF  L=)E"2=tD ?L K1 K+ ++  A %  Mrc[U[5(a[UR[ U55$[U[ [ 45(a[URU5$UR$r) isinstancerrxrrOrrrGr)layoutrs rrGrGsW&122V..E ;; F[*:; < <V]]E22###r)T) dataclassesrrtypingrtriton.language.corerrrtriton.runtime.jitr rYr r%r8rCrrrrrrrrr1r`rxrGrrrrs(TT1  S S $ 3" 3 3 $ 8' 8 8 $:%::z $'X#'X'XT $6/66r $0(00f $4 ('.('!('V$$& $[a [a[a| $4 .y<.y!.yb $4 yNyN!yNx $! ! ! J,$r