Pi0jddlZddlmZddlZddlmZerddlZddlmZ ddlm Z ddDZ ej e gdej d e jd e jd e jd e jd e jde jf dZejdd d)dejdejdejdejdedejf dZejd)dZej de jfdZdejfdejdedejdeejejffd Zej de jfd!Zdejfdejdedeejejffd"Zej de jfd#Z d)dejd$ejdedejfd%ZdS d)dejdejdejdejdedejf d&Zdejfdejdedejdeejejffd'Zdejfdejdedeejejffd(ZdS)*N)Tuple) has_tritonConfigc Lg|]!}dD]}dD]}t||d|d"S)) @)) BLOCK_SIZE_M BLOCK_SIZE_N) num_stages num_warpsr).0block_mblock_nrs y/home/jaya/work/projects/VOICE-AGENT/VIET/agent-env/lib/python3.11/site-packages/torchao/kernel/blockwise_quantization.py rsx    $   &    $g > >!        )rr r )NKM_BUCKET BLOCK_SIZE_K)configskeyrrrrrrc tjd} tjd} tj|| }| | ztjd| z|z}| | ztjd| z|z}tjd| }||dddf|zz|dddfz}||dddf|zz|dddfz}|||zz}||| z|zz}tj| | ftj}t |D]}tj||dddf||| zz kd}tj||dddf||| zz kd}tj|}tj|}|tj|||dddfz|dddfzz }|| z }|| z }|dz }|dz }| |j j }| | ztjd| z}| | ztjd| z}||dddf|zz|dddfz}|dddf|k|dddf|kz}tj |||dS)Nraxisdtypeg)maskotherr') tl program_idcdivarangezerosfloat32rangeloaddottor& element_tystore)a_ptrb_ptrc_ptra_s_ptrb_s_ptrMrrrrrrpid_mpid_nkoffs_moffs_noffs_ka_ptrsb_ptrsa_s_ptrsb_s_ptrs accumulatoriaba_sb_scc_ptrsr's rblockwise_fp8_gemm_kernelrNs$ 1%%% 1%%% GA| $ $,&1l)C)CCqH,&1l)C)CCqH1l++41,,vdAAAg>aaa1,,vaaag>VaZ'f 499h l;2:NNN q  AVD!!!G_q1|;K7K%KSVWWWAVAAAtG_q1|;K7K%KSVWWWA'(##C'(##C 26!Q<<#aaag,6T111WE EK l "F l "F MH MHH NN5;1 2 2% !\(B(BB% !\(B(BB41,,vdAAAg>qqq$w!#tQQQw!(;< &&&&&&rzao::blockwise_fp8_gemm) mutates_argsr rHrJrIrK block_sizereturnc > |sJ|sJ|sJ|sJ|d}||z |d tjtj }|jg|dd Rdtji} fd}t|||||| ||| |S)Nrr&cptj|dtj|dfS)Nrrtritonr,)METAr;rs rz$blockwise_fp8_gemm..`s2 K4/ 0 0 K4/ 0 0 r)r) is_contiguoussizenumelmathceillog2 new_emptytorchbfloat16rN) rHrJrIrKrQrrrLgridr;rs @@rblockwise_fp8_gemmrdOs3            """""  """"" FF2JJ GGIIN FF1II9TYq\\** AK @#2# @ @ @ @ @ @      "$' q!S#q!Qz    rc|d}|jg|dd|Rdtji}|S)NrrTr&)r[r`rarb)rHrJrIrKrQrrLs r_rfisL FF1II AK @#2# @ @ @ @ @ @r BLOCK_SIZEctjd}||ztjd|z}tj||ztj}tjtj|dz }||z }||jj }tj ||z|tj ||z|dS)a! Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factor in `s_ptr`. Args: x_ptr (triton.Pointer): Pointer to the input tensor. y_ptr (triton.Pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (triton.Pointer): Pointer to the output tensor where scaling factors will be stored. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. Returns: None rr"|@N) r*r+r-r1r3r/maxabsr&r4r5) x_ptry_ptrs_ptrrgpidoffsxsys rfp8_blockwise_act_quant_kernelrtosm###Z")Az":":: GEDL ! ! $ $RZ 0 0 F26!99   % E DD' ( ( q!!! a     rrqr&cs Jdd|zdksJd|d|tjtjfvs Jdtj|}jgddd|zRd tji}fd }t|||| ||fS) a Quantizes the input tensor `x` using block-wise quantization with block size being BLOCK_SIZEx1. Args: x (torch.Tensor): The input tensor to be quantized. Must be contiguous and its last dimension size must be divisible by `block_size`. block_size (int, optional): The size of the blocks to be used for quantization. Default is 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Default is `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. Input tensor must be contiguousrTrz@Last dimension size must be divisible by block_size (block_size=)6dtype must be torch.float8_e4m3fn or torch.float8_e5m2r%Nr&c`tj|dfSNrg)rWr,r\)metarqs rrYz)fp8_blockwise_act_quant..s#V[D4FGGIrrg) rZr[ra float8_e4m3fn float8_e5m2 empty_liker`r/rt)rqrQr&rsrrrcs` rfp8_blockwise_act_quantrs#"  CC"CCC vvbzzJ&!+++ \z \ \ \,++         D     Qe , , , AK V#2# Vr j(@ V V V V VIIII&t,Q1LLLL!t rctjd}tjd}tj||}||ztjd|z} ||ztjd|z} | dddf|z| dddfz} | dddf|k| dddf|kz} tj|| z| tj} tjtj| dz }| |z }||j j }tj || z|| tj |||zz|z|dS)aj Quantizes the input tensor `x_ptr` and stores the result in `y_ptr` and the scaling factors in `s_ptr`. Args: x_ptr (tl.pointer): Pointer to the input tensor. y_ptr (tl.pointer): Pointer to the output tensor where quantized values will be stored. s_ptr (tl.pointer): Pointer to the output tensor where scaling factors will be stored. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): The size of the block to be processed by each program instance. rr"r$Nr)ri) r*r+r,r-r1r3r/rjrkr&r4r5)rlrmrnr;rrgr<r=nr?r@rpr'rqrrrss r!fp8_blockwise_weight_quant_kernelrsf 1%%% 1%%% GAz " "#bi:&>&>>#bi:&>&>>aaag"VD!!!G_4qqq$w!#tQQQw!(;< GEDLt , , , / / ; ; F26!99   % E DD' ( ( qt,,,, "U*A.....rcN|s Jd|dks Jd|d|zdkr|d|zdksJd|d|tjtjfvs Jd|\tj|| }||z|ztj }fd }t||||| ||fS) a Quantizes the given weight tensor using block-wise quantization with block size being BLOCK_SIZExBLOCK_SIZE. Args: x (torch.Tensor): The weight tensor to be quantized. block_size (int, optional): The block size to use for quantization. Defaults to 128. dtype (torch.dtype, optional): The dtype to use for the quantized tensor. Defaults to `torch.float8_e4m3fn`. Returns: Tuple[torch.Tensor, torch.Tensor]: A tuple containing: - The quantized weight tensor with dtype `dtype`. - A tensor of scaling factors with dtype `torch.float32`. rvz#Input tensor must have 2 dimensionsrr$zABoth dimensions of x must be divisible by block_size (block_size=rwrxr%cptj|dtj|dfSrzrVr{r;rs rrYz,fp8_blockwise_weight_quant..2 K4 - . . K4 - . . rr|) rZdimr[rar}r~rr`r/r)rqrQr&rsrrrcr;rs @@rfp8_blockwise_weight_quantrs[   CC"CCC uuww!|||B|||vvayy:%**qvvayy:/E/J/J/J ]PZ ] ] ]0K/JJ         D   vvxx1  Qe , , , KKZj K N N      *$/1aA*UUUU!t rcNtjd}tjd}tj||}||ztjd|z} ||ztjd|z} | dddf|z| dddfz} | dddf|k| dddf|kz} tj|| z| tj} tj|||zz|z}| |z}tj|| z|| dS)a% Dequantizes weights using the provided scaling factors and stores the result. Args: x_ptr (tl.pointer): Pointer to the quantized weights. s_ptr (tl.pointer): Pointer to the scaling factors. y_ptr (tl.pointer): Pointer to the output buffer for dequantized weights. M (int): Number of rows in the weight matrix. N (int): Number of columns in the weight matrix. BLOCK_SIZE (tl.constexpr): Size of the block for tiling. Returns: None rr"r$Nr))r*r+r,r-r1r3r/r5)rlrnrmr;rrgr<r=rr?r@rpr'rqrrrss r#fp8_blockwise_weight_dequant_kernelrs1$ 1%%% 1%%% GAz " "#bi:&>&>>#bi:&>&>>aaag"VD!!!G_4qqq$w!#tQQQw!(;< GEDLt , , , / / ; ; GEEAI%- . . E qt,,,,,,rrrc|r|s Jd|dkr|dks Jd|\tj|tj}fd}t ||||||S)a' Dequantizes the given weight tensor using the provided scale tensor. Args: x (torch.Tensor): The quantized weight tensor of shape (M, N). s (torch.Tensor): The scale tensor of shape (M, N). block_size (int, optional): The block size to use for dequantization. Defaults to 128. Returns: torch.Tensor: The dequantized weight tensor of the same shape as `x`. Raises: AssertionError: If `x` or `s` are not contiguous or if their dimensions are not 2. z Input tensors must be contiguousrz$Input tensors must have 2 dimensionsr%cptj|dtj|dfSrzrVrs rrYz.fp8_blockwise_weight_dequant..rrr|)rZrr[rarget_default_dtyper)rqrrrQrsrcr;rs @@rfp8_blockwise_weight_dequantrs"   Q__%6%6  .  6uuww!||1 .T ,vvxx1  Qe&=&?&? @ @ @      ,D1!Q1aJWWWWrc tdNzunsupported without tritonAssertionError)rHrJrIrKrQs rrdrd#s9:::rc tdrrrqrQr&s rrr,9:::rc tdrrrs rrr1rr)r )r]typingrratorch.utils._tritonrrWtriton.languagelanguager*rfp8_gemm_configsautotunejit constexprrNlibrary custom_opTensorintrd register_fakerfrtr}r&rrrrrrOrrrsY  ****** :<<g;MMM         )   V_ &L&L&L Z*' <*' <*',*'l*'l*'l*'*'*'Z*'X ]5BGG   < \ <\     HG2%&%  Z! !!!Z!.,/UEX <%(7<{ u|U\) *> Z//1|///Z/:,/e6I!! <!%(! u|U\) *!!!!F Z-/1|---Z-<=@ <!L69 J ;; <; \; <;\ ;  ;  ;;;;,/UEX;; <;%(;7<{; u|U\) *;;;; ,/e6I;; <;%(; u|U\) *;;;;;;r