hWE SSKJrJr SSKJrJr SSKr\"SS9"SS55r\"5rS\ S \ 4S jr \R4S \RS \ S \ S\RS\ S \R4 Sjjr \R4S\RS \ S \ S\RS\ S \R4 SjjrS\ S \\R\\ \\ 44SjrS\ S \\R\R\R44SjrSrSrSrg)) dataclassfield)ListTupleNT)frozenc|\rSrSr%Sr\\S'Sr\\S'Sr\\S'\ "SS 9r \ \\S '\ "S S 9r \ \\S 'S r g)Marlin24Constants TILE MIN_THREAD_N@ MAX_PARALLELc SS/$)NrW/home/james-whalen/.local/lib/python3.13/site-packages/torchao/sparsity/marlin/utils.pyMarlin24Constants.s1a&r)default_factorySUPPORTED_NUM_BITSc /SQ$)N) rr rrrrrrsEVrSUPPORTED_GROUP_SIZESrN)__name__ __module__ __qualname____firstlineno__r int__annotations__rrrrrr__static_attributes__rrrr r sJD#NL#L#%*.$IS I',=V'W49Wrr num_bitsreturncHU[R;d SU35eSU-$)zCompute the packing factor for a given number of bits. Args: num_bits (int): Number of bits to pack. Returns: int: The packing factor. zUnsupported num_bits = r)constr)r&s rget_pack_factorr*s/ u// /U3J8*1UU / >rq_wsize_ksize_npermtilecURX4:XdeX-S:Xd SUSU35eX$-S:Xd SUSU35eURX-XBU-U45nURS5nURX-X$-45nURSUR545SS2U4RUR5nU$)avPermute weights to 16x64 Marlin tiles. Args: q_w (torch.Tensor): Quantized weights. size_k (int): Number of input features. size_n (int): Number of output features. perm (torch.Tensor): The computed permutation tensor to be applied. tile (int, optional): Tile size. Defaults to `TILE`. Returns: torch.Tensor: Weight tensor permuted to Marlin tiles. r size_k = , tile = rrN)shapereshapepermutenumel)r+r,r-r.r/s rmarlin_permute_weightsr;'s& 99( (( ( =A B6()D6BB  =A B6()D6BB  ++v~tt^TB CC ++l #C ++v~v}5 6C ++r4::<( )!T' 2 : :399 EC Jr q_w_unpacked reverse_permcURSU4X-URSU-4:XdeX-S:Xd SUSU35eX$-S:Xd SUSU35eURSUR545SS2U4RUR5nURX-X$-XD45nURS5nURX45nU$)aReverse permute weights from 16x64 Marlin tiles. Args: q_w_unpacked (torch.Tensor): Unpacked quantized weights. size_k (int): Number of input features. size_n (int): Number of output features. reverse_perm (torch.Tensor): The computed reverse permutation tensor to be applied. tile (int, optional): Tile size. Defaults to `TILE`. Returns: torch.Tensor: Weight tensor reverse permuted from Marlin tiles. rr5r1r2rNr3)r7r8r:r9)r<r,r-r=r/q_w_comps rreverse_marlin_permute_weightsr@Hs $   q !6 *1%/   =A B6()D6BB  =A B6()D6BB ##R););)=$>? < gl  !  LMH -H 01H Orc /n[S5Hn/nUS-nUS-nSH]nSUS--SUS--S-SUS-S--SUS-S--S-4H,nURSU-US--SUS---SU--5 M. M_ [S5H*nURUV s/sH oSU--PM sn 5 M, M [R"U[R S 9n US:Xa&[R"/S Q[R S 9n OFUS:Xa&[R"/S Q[R S 9n O[ S RU55eU RS [U 55SS2U 4RS 5n /n [S5H*nU RSVs/sH oS-U-PM sn5 M, /n [S5H+nU RSVs/sH nSU-U-PM sn5 M- XU 4$s sn fs snfs snf)aPrecompute permutations for Marlin24 weight and scale shuffling Marlin works on [16*2,64] tiles. The goal of the permutations is to reorder the weight data so that it is compatible with the tensor-core format that is described here: https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#matrix-fragments-for-mma-m16n8k16-with-floating-point-type As a result of this reordering, the vector loads inside the kernel will get the data as it is needed for tensor-core (without the need to use ldmatrix instructions) Args: num_bits (int): Number of bits to pack. Returns: Tuple[torch.Tensor, List[int], List[int]]: The weight permutation tensor, scale permutation list, and scale permutation list for a single group. rrr4)rr5r5r r)dtype)rr4rr5r6r3znum_bits must be 4 or 8, got {}rN)rrr5rEr4rDr6rF)rr5r4r6rrErDrF) rangeappendextendtorchtensorint32 ValueErrorformatviewlenr8)r& perm_listiperm1colcol_oblockrowjpr. interleave scale_permscale_perm_singles r get_perms_24r]ls  I 2Y1fqEQU QU aQUQYQUQY!#   R#X 3a37mCa%iOP qA   7A!a%i7 8 <<  5D1}\\":%++N Q\\,ekkB :AA(KLL 99RZ )!Z- 8 @ @ DDJ 1X.FG.Fq519.FGH$& 1X  5M!N5M!a%!)5M!NO . ../8"H"OsG> "H H c[U5upnUR5n[R"U5R5n[R"U5R5nXEU4$)a2Reverse permutation for Marlin24 weight and scale shuffling from `get_perms_24`. Args: num_bits (int): Number of bits to pack. Returns: Tuple[torch.Tensor, torch.Tensor, torch.Tensor]: The reversed weight permutation tensor, scale permutation list and scale permutation list for single group. )r]argsortrJrK)r&perm_24 scale_perm_24scale_perm_single_24r.r[r\s rget_reverse_perms_24rcs]4@3I0G0 ?? Dm,446J %9:BBD . ..rcX[R"SXS9SS2S4RSU5n[R"SXS9RUS5nSnURS:XaSOSnXF-U-US-S--US-S --XG-S -S-S--XF-S-S --nUS-S:HUS-S:H-R [R 5nUS-S:HUS-S:H-R [R 5n XHU - - nXXU - -nSn XZ-n XZ-n X-U -XJ--U -R S 5$) Nrdevicer5rr4rr rrr)rJarangerepeatitemsizetoint8rO) m meta_ncols meta_dtyperfdst_rowsdst_colsgroup_xgroup_ytopright bottomleftrZcols_majcols_mins r*_calculate_meta_reordering_scatter_offsetsrwsu||Aq0D9@@JOH||Az9@@AFHG''1,b"G g% a<1   a zL <<++B/L A F Fq UD[[ 223F "+q!Qw19,q!Qw19,q!Qw19,q!Qw19,q!Qw2:-q!Qw2:-q!Qw2:-q!Qw 1 $+q!Qw19,q!Qw19,q!Qw19,q!Qw19,q!Qw2:-q!Qw2:-q!Qw2:-q!Qw2:-q!Qw2:-q!Qw BJ$.q!Rx BJ$.q!Rx BJ$.q!Rx BJ$.q!Rx BJ$.q!Rx BJ$.q!RxKKO QA W,V  ::aQ r) dataclassesrrtypingrrrJr r)r#r*r Tensorr;r@r]rcrwrrrrrrsb )  $XXX  c c $     ,,     \\ L !,,! ! !,, !  !  \\ !H5/35/5tCy$s))K#L5/p// 5<<u|| 34/BSB}8H^ r