PiWE ddlmZmZddlmZmZddlZedGddZeZde d e fd Z ej fd ej d e d e dej de d ej f dZ ej fdej d e d e dej de d ej f dZde d eej ee ee ffdZde d eej ej ej ffdZdZdZdZdS)) dataclassfield)ListTupleNT)frozenceZdZUdZeed<dZeed<dZeed<edZ e eed <ed Z e eed <d S) Marlin24ConstantsTILE MIN_THREAD_N@ MAX_PARALLELc ddgS)Nrq/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/torchao/sparsity/marlin/utils.pyzMarlin24Constants.s 1a&r)default_factorySUPPORTED_NUM_BITSc gdS)N) rr rrrrrzMarlin24Constants.s EVEVEVrSUPPORTED_GROUP_SIZESN) __name__ __module__ __qualname__r int__annotations__r rrrrrrrrr r sD#NNNL#L#%*E..$I$I$IS III',u=V=V'W'W'W49WWWWWrr num_bitsreturncB|tjvs Jd|d|zS)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&s3 u/ / / /1U81U1U / / / >rq_wsize_ksize_npermtilec|j||fksJ||zdksJd|d|||zdksJd|d||||z|||z|f}|d}|||z||zf}|d|fdd|f|j}|S)aPermute 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_weightsr7's & 9( ( ( ( ( D=A   B6BBDBB    D=A   B6BBDBB    ++v~tVt^TB C CC ++l # #C ++v~v}5 6 6C ++r4::<<( ) )!!!T' 2 : :39 E EC Jr q_w_unpacked reverse_permc|jd|f||z|jd|zfksJ||zdksJd|d|||zdksJd|d||d|fdd|f|j}|||z||z||f}|d}|||f}|S)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. rr1r-r.rNr/)r3r4r6r5)r8r(r)r9r+q_w_comps rreverse_marlin_permute_weightsr<HsA$  q !6 *$1%/     D=A   B6BBDBB    D=A   B6BBDBB   ##R););)=)=$>?? < gl !! 44t LMMH --H 011H Orc  g}tdD] g} dz}|dz}dD]\}d dzzd dzzdzd dzdzzd dzdzzdzfD]/}|d|z|dzzd|dzzzd|zz0]tdD]# | fd |D$tj|tj }|dkr#tjgd tj }nK|dkr#tjgd tj }n"t d ||dt|dd|f d}g} tdD]# |  fddD$g} tdD]# |  fddD$|| | fS)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. rrr0)rr1r1r rc g|] }|dzz S)r1r).0pjs r z get_perms_24..s!777Aa!a%i777r)dtype)rr0rr1r2r/znum_bits must be 4 or 8, got {}rNc g|] }dz|z Srrr@rBis rrCz get_perms_24..s!GGG1q519GGGr)rrr1rFr0rEr2rGc g|] }dz|z SrIrrJs rrCz get_perms_24..s!!N!N!N!a%!)!N!N!Nr)rr1r0r2rrFrErG) rangeappendextendtorchtensorint32 ValueErrorformatviewlenr4) r" perm_listperm1colcol_oblockrowr* interleave scale_permscale_perm_singlerKrBs @@r get_perms_24r`lsy I 2YY 9 91fq Q QEQU QU aQUQYQUQY!#  Q Q  R#X 3a37mCa%iOPPPP  Qq 9 9A   7777777 8 8 8 8 9 <  5 5 5D1}}\":":":%+NNN Q\,,,ekBBB :AA(KKLLL 99RZ ) )!!!Z- 8 @ @ D DDJ 1XXIIGGGG.FGGGHHHH#% 1XXPP  !N!N!N!N5M!N!N!NOOOO . ..rct|\}}}|}tj|}tj|}|||fS)aJReverse 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`argsortrPrQ)r"perm_24 scale_perm_24scale_perm_single_24r*r^r_s rget_reverse_perms_24rfso4@3I3I0G]0 ??  Dm,,4466J %9::BBDD . ..rctjd||dddfd|}tjd|||d}d}|jdkrdnd}||z|z|dzdzz|dzd zz||zd zdzdzz||zdzd zz}|dzdk|dzdkztj}|dzdk|dzdkztj} ||| z z }||| z z}d} || z} || z} | |z| z|| zz| zd S) Nrdevicer1rr0rr rrr)rParangerepeatitemsizetoint8rU) m meta_ncols meta_dtyperidst_rowsdst_colsgroup_xgroup_ytopright bottomleftr]cols_majcols_mins r*_calculate_meta_reordering_scatter_offsetsrzs|Aq000D9@@JOOH|Az&999@@AFFHG'1,,bb"G Gg% a<1   a :z6L <++B//L A A F Fq* U UD[ J223F ""+qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw 1 $ $+qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw19,qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw2:-qqq!!!Qw BJ$.qqq!!!Rx BJ$.qqq!!!Rx BJ$.qqq!!!Rx BJ$.qqq!!!Rx BJ$.qqq!!!Rx BJ$.qqq!!!RxKKOO QA W,V<<>    ::aQ  r) dataclassesrrtypingrrrPr r%r r&r Tensorr7r<r`rfrzrrrrrrs. )(((((((  $XXXXXXXX  c c    $     ,     \ L !!,! ! !, !  !  \ !!!!H5/35/5tCy$s))K#L5/5/5/5/p// 5<u| 34////BSSSB}8}8}8H^ ^ ^ ^ ^ r